Ë
    ÄÓ^jù ã                  óÌ  — d dl mZ d dlZd dlZd dlZd dlZd dlZd dlmZm	Z	m
Z
 d dlZd dlZd dlmZ d dlmZmZmZ ddlmZ ddlmZ dd	lmZmZ dd
lmZ ddlmZ ddl m!Z!m"Z"m#Z#m$Z$m%Z% ddl&m'Z'm(Z(m)Z)  G d„ de%«      Z* e*«       jV                  Z,e
rd dl-m.Z.m/Z/ ddl0m1Z1 ddl2m3Z3 ddl4m5Z5 dZ6ejn                  jq                  e9d«      Z: G d„ d«      Z; G d„ de<«      Z= G d„ de$«      Z>ej~                   G d„ d«      «       Z@ej~                   G d„ d «      «       ZAej~                   G d!„ d"«      «       ZBej~                   G d#„ d$«      «       ZC G d%„ d&e(«      ZD G d'„ d(e)«      ZEy))é    )ÚannotationsN)ÚAnyÚcastÚTYPE_CHECKING)Ú
OrderedSet)ÚMaxÚMinÚModularIndexingé   )Úconfig©Útorch_dtype_to_jax)Úget_fused_kernel_nameÚget_kernel_metadata)ÚVé   )ÚBlockPatternMatcher)ÚBackendFeatureÚCSEVariableÚIndentedBufferÚOpOverridesÚPythonPrinter)ÚIterationRangesEntryÚ
SIMDKernelÚSIMDSchedulingc                  ó(   — e Zd ZdZdd„Zdd„Zdd„Zy)ÚPallasPrinterzO
    Custom sympy printer for Pallas that handles JAX-specific constructs.
    c                óÎ   — | j                  |j                  d   «      }| j                  |j                  d   «      }| j                  |j                  d   «      }d|› d|› d|› d�S )z!Convert sympy Where to jnp.where.r   r   r   ú
jnp.where(ú, ú))ÚdoprintÚargs)ÚselfÚexprÚcÚpÚqs        úi/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/_inductor/codegen/pallas.pyÚ_print_WherezPallasPrinter._print_Where(   s_   € à�L‰L˜Ÿ™ 1™Ó&ˆØ�L‰L˜Ÿ™ 1™Ó&ˆØ�L‰L˜Ÿ™ 1™Ó&ˆØ˜A˜3˜b   2 a S¨Ð*Ð*ó    c                ó�   — |j                   D �cg c]  }| j                  |«      ‘Œ }}|d   }|dd D ]  }d|› d|› d�}Œ |S c c}w )z7Convert sympy Min to jnp.minimum for JAX compatibility.r   r   Nújnp.minimum(r    r!   ©r#   r"   ©r$   r%   Úargr#   Úresults        r)   Ú
_print_MinzPallasPrinter._print_Min/   ó^   € à-1¯Y©YÖ7 c�—‘˜SÕ!Ð7ˆÐ7Ø�a‘ˆØ˜˜�8ò 	5ˆCØ# F 8¨2¨c¨U°!Ð4‰Fð	5àˆùò	 8ó   �Ac                ó�   — |j                   D �cg c]  }| j                  |«      ‘Œ }}|d   }|dd D ]  }d|› d|› d�}Œ |S c c}w )z7Convert sympy Max to jnp.maximum for JAX compatibility.r   r   Nújnp.maximum(r    r!   r.   r/   s        r)   Ú
_print_MaxzPallasPrinter._print_Max7   r3   r4   N)r%   ú
sympy.ExprÚreturnÚstr)Ú__name__Ú
__module__Ú__qualname__Ú__doc__r*   r2   r7   © r+   r)   r   r   #   s   „ ñó+óôr+   r   )ÚCallableÚSequence)ÚIRNode)ÚReductionType)ÚBaseSchedulerNodeÚmainÚkernel_codec                  ó&   — e Zd ZdZddd„Zddœd„Zy)ÚPallasKernelWrapperz6Wrapper to provide .run() interface for Pallas kernelsNc                óL   — || _         || _        t        j                  d|«       y )NzPallas kernel path: %s)Ú	kernel_fnÚkernel_pathÚkernel_code_logÚinfo)r$   rJ   rK   s      r)   Ú__init__zPallasKernelWrapper.__init__V   s"   € Ø"ˆŒØ&ˆÔÜ×ÑÐ5°{ÕCr+   )Ústreamc               ó*   —  | j                   |d|i|¤ŽS )a5  
        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
        rO   )rJ   )r$   rO   r#   Úkwargss       r)   ÚrunzPallasKernelWrapper.run[   s   € ð ˆt�~‰~˜tÐ=¨FÐ=°fÑ=Ð=r+   ©N)rJ   zCallable[..., Any]rK   ú
str | None)r;   r<   r=   r>   rN   rR   r?   r+   r)   rH   rH   S   s   „ Ù@ôDð
 !%õ >r+   rH   c                  ó   — e Zd ZdZy)ÚUnsupportedzJException raised when an operation is not supported by the Pallas backend.N)r;   r<   r=   r>   r?   r+   r)   rV   rV   j   s   „ ÚTr+   rV   c                  ó.  — e Zd ZdZedtd„«       Zedtd„«       Zedtd„«       Zedtd„«       Zedtd„«       Z	edtd„«       Z
edtd„«       Zedtd	„«       Zedtd
„«       Zedtd„«       Zedtd„«       Zedtd„«       Zedtd„«       Zedtd„«       Zedtd„«       Zedtd„«       Zedtd„«       Zedtd„«       Zedtd„«       Zedtd„«       Zedtd„«       Zedtd„«       Zedtd„«       Zedtd„«       Zedtd„«       Zedtd„«       Zedud„«       Zedud„«       Z edud„«       Z!edvd„«       Z"edwd „«       Z#e	 	 dx	 	 	 	 	 	 	 	 	 dyd"„«       Z$edzd#„«       Z%ed{d$„«       Z&ed{d%„«       Z'ed|d&„«       Z(edtd'„«       Z)edtd(„«       Z*edtd)„«       Z+edtd*„«       Z,edtd+„«       Z-edtd,„«       Z.edud-„«       Z/edud.„«       Z0edud/„«       Z1edud0„«       Z2edud1„«       Z3edtd2„«       Z4edtd3„«       Z5edtd4„«       Z6edud5„«       Z7edud6„«       Z8edud7„«       Z9edtd8„«       Z:edud9„«       Z;edud:„«       Z<edud;„«       Z=edud<„«       Z>edud=„«       Z?edud>„«       Z@edud?„«       ZAed}d@„«       ZBeBZCedtdA„«       ZDedtdB„«       ZEedtdC„«       ZFedtdD„«       ZGedtdE„«       ZHedtdF„«       ZIedtdG„«       ZJedtdH„«       ZKedtdI„«       ZLedtdJ„«       ZMedtdK„«       ZNedtdL„«       ZOeMZPedtdM„«       ZQeNZRedtdN„«       ZSed~dO„«       ZTed~dP„«       ZUeTZVeUZWed~dQ„«       ZXedtdR„«       ZYed~dS„«       ZZed~dT„«       Z[ed~dU„«       Z\eddV„«       Z]eddW„«       Z^eddX„«       Z_eddY„«       Z`eddZ„«       Zaedd[„«       Zbedd\„«       Zcedd]„«       Zdedd^„«       Zeedd_„«       Zfedd`„«       Zgedda„«       Zhedtdb„«       Ziedtdc„«       Zjed€dd„«       Zkedude„«       Zledudf„«       Zmedudg„«       Znedtdh„«       Zoedtdi„«       Zpedudj„«       Zqedudk„«       Zredudl„«       Zsedtdm„«       Ztedudn„«       Zuedudo„«       Zved�dp„«       Zwed‚dq„«       Zxed‚dr„«       Zyedƒds„«       Zzy!)„ÚPallasKernelOverrideszª
    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                ó   — d| › d�S )Nzjnp.sin(r!   r?   ©Úxs    r)   ÚsinzPallasKernelOverrides.sinv   ó   € ð ˜!˜˜AˆÐr+   c                ó   — d| › d�S )Nzjnp.cos(r!   r?   rZ   s    r)   ÚcoszPallasKernelOverrides.cos{   r]   r+   c                ó   — d| › d�S )Nzjnp.tan(r!   r?   rZ   s    r)   ÚtanzPallasKernelOverrides.tan€   r]   r+   c                ó   — d| › d�S )Nz	jnp.sinh(r!   r?   rZ   s    r)   ÚsinhzPallasKernelOverrides.sinh…   ó   € ð ˜1˜#˜QÐÐr+   c                ó   — d| › d�S )Nz	jnp.cosh(r!   r?   rZ   s    r)   ÚcoshzPallasKernelOverrides.coshŠ   rd   r+   c                ó   — d| › d�S )Nz	jnp.tanh(r!   r?   rZ   s    r)   ÚtanhzPallasKernelOverrides.tanh�   rd   r+   c                ó   — d| › d�S )Nzjnp.arcsin(r!   r?   rZ   s    r)   ÚasinzPallasKernelOverrides.asin”   ó   € ð ˜Q˜C˜qÐ!Ð!r+   c                ó   — d| › d�S )Nzjnp.arccos(r!   r?   rZ   s    r)   ÚacoszPallasKernelOverrides.acos™   rk   r+   c                ó   — d| › d�S )Nzjnp.arctan(r!   r?   rZ   s    r)   ÚatanzPallasKernelOverrides.atanž   rk   r+   c                ó   — d| › d�S )Nzjnp.exp(r!   r?   rZ   s    r)   ÚexpzPallasKernelOverrides.exp£   r]   r+   c                ó   — d| › d�S )Nz	jnp.exp2(r!   r?   rZ   s    r)   Úexp2zPallasKernelOverrides.exp2¨   ó   € à˜1˜#˜QÐÐr+   c                ó   — d| › d�S )Nz
jnp.expm1(r!   r?   rZ   s    r)   Úexpm1zPallasKernelOverrides.expm1¬   ó   € à˜A˜3˜aÐ Ð r+   c                ó   — d| › d�S )Nzjnp.log(r!   r?   rZ   s    r)   ÚlogzPallasKernelOverrides.log°   r]   r+   c                ó   — d| › d�S )Nz
jnp.log10(r!   r?   rZ   s    r)   Úlog10zPallasKernelOverrides.log10µ   rw   r+   c                ó   — d| › d�S )Nz	jnp.log2(r!   r?   rZ   s    r)   Úlog2zPallasKernelOverrides.log2¹   rt   r+   c                ó   — d| › d�S )Nz
jnp.log1p(r!   r?   rZ   s    r)   Úlog1pzPallasKernelOverrides.log1p½   rw   r+   c                ó   — d| › d�S )Nz	jnp.sqrt(r!   r?   rZ   s    r)   ÚsqrtzPallasKernelOverrides.sqrtÁ   rd   r+   c                ó   — d| › d�S )Nzjax.lax.rsqrt(r!   r?   rZ   s    r)   ÚrsqrtzPallasKernelOverrides.rsqrtÆ   s   € ð   ˜s !Ð$Ð$r+   c                ó   — d| › d�S )Nzjnp.abs(r!   r?   rZ   s    r)   ÚabszPallasKernelOverrides.absË   r]   r+   c                ó   — d| › d�S )Nz(-r!   r?   rZ   s    r)   ÚnegzPallasKernelOverrides.negÐ   s   € à�A�3�aˆyÐr+   c                ó   — d| › d�S )Nz
jnp.floor(r!   r?   rZ   s    r)   ÚfloorzPallasKernelOverrides.floorÔ   ó   € ð ˜A˜3˜aÐ Ð r+   c                ó   — d| › d�S )Nz	jnp.ceil(r!   r?   rZ   s    r)   ÚceilzPallasKernelOverrides.ceilÙ   rd   r+   c                ó   — d| › d�S )Nz
jnp.trunc(r!   r?   rZ   s    r)   ÚtrunczPallasKernelOverrides.truncÞ   rŠ   r+   c                ó   — d| › d�S )Nz
jnp.round(r!   r?   rZ   s    r)   ÚroundzPallasKernelOverrides.roundã   rŠ   r+   c                ó   — d| › d�S )Nzjax.nn.sigmoid(r!   r?   rZ   s    r)   ÚsigmoidzPallasKernelOverrides.sigmoidè   ó   € à    1Ð%Ð%r+   c                ó   — d| › d�S )Nr6   z, 0)r?   rZ   s    r)   ÚreluzPallasKernelOverrides.reluì   s   € à˜a˜S Ð%Ð%r+   c                ó   — d| › d|› d�S )Nz
jnp.power(r    r!   r?   ©ÚaÚbs     r)   ÚpowzPallasKernelOverrides.powð   s   € à˜A˜3˜b   1Ð%Ð%r+   c                ó   — d| › d|› d�S )Nr6   r    r!   r?   r—   s     r)   ÚmaximumzPallasKernelOverrides.maximumô   ó   € ð ˜a˜S  1 # QÐ'Ð'r+   c                ó   — d| › d|› d�S )Nr-   r    r!   r?   r—   s     r)   ÚminimumzPallasKernelOverrides.minimumù   r�   r+   c                ó   — d| › d|› d|› d�S )Nr   r    r!   r?   )Úcondr˜   r™   s      r)   ÚwherezPallasKernelOverrides.whereþ   s   € ð ˜D˜6  A 3 b¨¨¨1Ð-Ð-r+   c                óä   —  |«       }t        |t        «      rCt        j                  |«      rd}n6t        j                  |«      r
|dkD  rdnd}nt        |«      }nt        |«      }d| › d|› d|› d�S )zŠ
        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Úotherr1   Ú	other_strs        r)   ÚmaskedzPallasKernelOverrides.masked  sl   € ñ “ˆä�eœUÔ#Ü�z‰z˜%Ô Ø%‘	Ü—‘˜EÔ"Ø).°ª™I¸
‘	ä  ›K‘	ä˜U›ˆIà˜D˜6  F 8¨2¨i¨[¸Ð:Ð:r+   Nc                óÆ   — |t         j                  k(  r;t        j                  j	                  «       j
                  dk(  rt         j                  }t        |«      }d| › d|› d�S )NÚtpuzjnp.asarray(ú	).astype(r!   )ÚtorchÚint64r   ÚgraphÚget_current_device_or_throwÚtypeÚint32r   )r[   ÚdtypeÚ	src_dtypeÚuse_compute_typesÚ	jax_dtypes        r)   Úto_dtypezPallasKernelOverrides.to_dtype  sS   € ð ”E—K‘KÒ¤A§G¡G×$GÑ$GÓ$I×$NÑ$NÐRWÒ$WÜ—K‘KˆEÜ& uÓ-ˆ	à˜a˜S 	¨)¨°AÐ6Ð6r+   c                óF   — t        |«      }t        |«      }d| › d|› d|› d�S )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   )r[   r»   r¼   r¾   Újax_src_dtypes        r)   Úto_dtype_bitcastz&PallasKernelOverrides.to_dtype_bitcast%  s7   € ô ' uÓ-ˆ	Ü*¨9Ó5ˆà:¸1¸#¸YÀ}ÀoÐUXÐYbÐXcÐcdÐeÐer+   c                ó  — ddl m} t        j                  j                  j                  t        j                  j                  | «      «       t        j                  j                  | «      }t        j                  j                  |«      }t        j                  j                  |«      }t        j                  j                  j                  t        j                  j                  | || «      ¬«      }t        j                  ||«      S )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ÚcomputerX   r¿   )r%   r»   rÄ   ÚpreparedÚrenamedÚidx_strÚvars          r)   Ú
index_exprz 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°5Ó9Ð9r+   c                ó.   — t         j                  | |«      S rS   )rX   rÕ   )r%   r»   s     r)   Ú
value_exprz PallasKernelOverrides.value_expr>  s   € ô %×/Ñ/°°eÓ<Ð<r+   c                óä   — t        |«      }|t        j                  k(  r| rdS dS t        | t        «      r4t        j                  | «      ryt        j                  | «      r	| dkD  rdS dS d| › d|› d	�S )
z/Convert a constant value to JAX representation.ÚTrueÚFalser¤   r   r¥   r¦   z
jnp.array(z, dtype=r!   )r   rµ   Úboolr§   r¨   r©   rª   r«   )Úvalr»   r¾   s      r)   ÚconstantzPallasKernelOverrides.constantD  sp   € ô ' uÓ-ˆ	Ø”E—J‘JÒÙ �6Ð- gÐ-ä�cœ5Ô!Ü�z‰z˜#ŒØ Ü�z‰z˜#ŒØ$'¨!¢G�yÐ;°Ð;Ø˜C˜5 ¨¨°1Ð5Ð5r+   c                ó   — d| › d�S )Nz	jnp.real(r!   r?   rZ   s    r)   ÚrealzPallasKernelOverrides.realR  rt   r+   c                ó   — d| › d�S )Nz	jnp.imag(r!   r?   rZ   s    r)   ÚimagzPallasKernelOverrides.imagV  rt   r+   c                ó   — d| › d�S )Nz	jnp.conj(r!   r?   rZ   s    r)   ÚconjzPallasKernelOverrides.conjZ  rt   r+   c                ó   — d| › d�S )Nz
jnp.angle(r!   r?   rZ   s    r)   ÚanglezPallasKernelOverrides.angle^  rw   r+   c                ó   — d| › d| › d�S )z8View complex tensor as real tensor with extra dimension.zjnp.stack([jnp.real(z), jnp.imag(z)], axis=-1)r?   rZ   s    r)   Úview_as_realz"PallasKernelOverrides.view_as_realb  s   € ð & a S¨°Q°C°|ÐDÐDr+   c                ó   — d| › d| › d�S )z#View real tensor as complex tensor.ú(z[..., 0] + 1j * z	[..., 1])r?   rZ   s    r)   Úview_as_complexz%PallasKernelOverrides.view_as_complexg  s   € ð �1�#Ð% a S¨	Ð2Ð2r+   c                ó   — d| › d|› d�S )Nré   z == r!   r?   r—   s     r)   ÚeqzPallasKernelOverrides.eqm  ó   € à�1�#�T˜!˜˜AˆÐr+   c                ó   — d| › d|› d�S )Nré   z != r!   r?   r—   s     r)   ÚnezPallasKernelOverrides.neq  rí   r+   c                ó   — d| › d|› d�S )Nré   z < r!   r?   r—   s     r)   ÚltzPallasKernelOverrides.ltu  ó   € à�1�#�S˜˜˜1ˆ~Ðr+   c                ó   — d| › d|› d�S )Nré   z <= r!   r?   r—   s     r)   ÚlezPallasKernelOverrides.ley  rí   r+   c                ó   — d| › d|› d�S )Nré   z > r!   r?   r—   s     r)   ÚgtzPallasKernelOverrides.gt}  rò   r+   c                ó   — d| › d�S )Nz
jnp.isnan(r!   r?   rZ   s    r)   rª   zPallasKernelOverrides.isnan�  rŠ   r+   c                ó   — d| › d�S )Nz
jnp.isinf(r!   r?   rZ   s    r)   r«   zPallasKernelOverrides.isinf†  rŠ   r+   c                ó   — d| › d�S )Nzjnp.isfinite(r!   r?   rZ   s    r)   ÚisfinitezPallasKernelOverrides.isfinite‹  s   € à˜q˜c Ð#Ð#r+   c                ó   — d| › d|› d�S )Nré   z >= r!   r?   r—   s     r)   ÚgezPallasKernelOverrides.ge�  rí   r+   c                ó   — d| › d|› d�S )Nzjnp.logical_and(r    r!   r?   r—   s     r)   Úlogical_andz!PallasKernelOverrides.logical_and”  ó   € ð " !  B q c¨Ð+Ð+r+   c                ó   — d| › d|› d�S )Nzjnp.logical_or(r    r!   r?   r—   s     r)   Ú
logical_orz PallasKernelOverrides.logical_or™  s   € ð !   2 a S¨Ð*Ð*r+   c                ó   — d| › d�S )Nzjnp.logical_not(r!   r?   rZ   s    r)   Úlogical_notz!PallasKernelOverrides.logical_notž  ó   € à! !  AÐ&Ð&r+   c                ó   — d| › d|› d�S )Nzjnp.logical_xor(r    r!   r?   r—   s     r)   Úlogical_xorz!PallasKernelOverrides.logical_xor¢  rÿ   r+   c                ó   — d| › d|› d�S )Nzjnp.arctan2(r    r!   r?   r—   s     r)   Úatan2zPallasKernelOverrides.atan2¨  r�   r+   c                ó   — d| › d|› d�S )Nz
jnp.hypot(r    r!   r?   r—   s     r)   ÚhypotzPallasKernelOverrides.hypot­  ó   € ð ˜A˜3˜b   1Ð%Ð%r+   c                ó   — d| › d|› d�S )Nz	jnp.fmod(r    r!   r?   r—   s     r)   ÚfmodzPallasKernelOverrides.fmod²  s   € ð ˜1˜#˜R ˜s !Ð$Ð$r+   c                ó   — d| › d|› d�S )Nzjnp.remainder(r    r!   r?   r—   s     r)   Ú	remainderzPallasKernelOverrides.remainder·  s   € à ˜s " Q C qÐ)Ð)r+   c                ó&   — d| › d|› d| › d|› d| › d�S )Nz
(jnp.sign(z) * jnp.sign(z) * (jnp.abs(z) // jnp.abs(z))).astype(ú.dtype)r?   r—   s     r)   ÚtruncdivzPallasKernelOverrides.truncdiv»  s.   € ð
 ˜A˜3˜m¨A¨3¨m¸A¸3¸mÈAÈ3ÈkÐZ[ÐY\Ð\cÐdÐdr+   c                ó   — d| › d|› d�S )Nré   z // r!   r?   r—   s     r)   ÚfloordivzPallasKernelOverrides.floordivÂ  rí   r+   c                ó   — d| › d|› d|› d�S )Nz	jnp.clip(r    r!   r?   )r[   Úmin_valÚmax_vals      r)   ÚclampzPallasKernelOverrides.clampÆ  s   € à˜1˜#˜R ˜y¨¨7¨)°1Ð5Ð5r+   c                ó   — d| › d| › d�S )Nzjnp.where(jnp.isnan(z), 0.0, jnp.sign(ú))r?   rZ   s    r)   ÚsignzPallasKernelOverrides.signÍ  s   € ð & a SÐ(9¸!¸¸BÐ?Ð?r+   c                ó   — d| › d�S )Nzjnp.signbit(r!   r?   rZ   s    r)   ÚsignbitzPallasKernelOverrides.signbitÓ  s   € ð ˜a˜S Ð"Ð"r+   c                ó   — d| › d�S )Nzjax.scipy.special.erf(r!   r?   rZ   s    r)   ÚerfzPallasKernelOverrides.erfÙ  s   € ð (¨ s¨!Ð,Ð,r+   c                ó   — d| › d�S )Nzjax.scipy.special.erfc(r!   r?   rZ   s    r)   ÚerfczPallasKernelOverrides.erfcÞ  s   € à(¨¨¨1Ð-Ð-r+   c                ó   — d| › d�S )Nzjax.scipy.special.erfinv(r!   r?   rZ   s    r)   ÚerfinvzPallasKernelOverrides.erfinvâ  s   € ð +¨1¨#¨QÐ/Ð/r+   c                ó   — d| › d�S )Nzjax.scipy.special.gammaln(r!   r?   rZ   s    r)   ÚlgammazPallasKernelOverrides.lgammaç  s   € ð ,¨A¨3¨aÐ0Ð0r+   c                ó   — d| › d�S )Nzjax.scipy.special.digamma(r!   r?   rZ   s    r)   ÚdigammazPallasKernelOverrides.digammaì  s   € à+¨A¨3¨aÐ0Ð0r+   c                ó   — d| › d| › d| › d�S )Nr   z>.astype(jnp.float64) == 0.0, 1.0, jax.scipy.special.bessel_jn(z&.astype(jnp.float64), v=0)[0]).astype(r  r?   rZ   s    r)   Ú	bessel_j0zPallasKernelOverrides.bessel_j0ð  ó+   € ð ˜˜ð +Ø+,¨#ð .Ø�c˜ð"ð	
r+   c                ó   — d| › d| › d| › d�S )Nr   z>.astype(jnp.float64) == 0.0, 0.0, jax.scipy.special.bessel_jn(z&.astype(jnp.float64), v=1)[1]).astype(r  r?   rZ   s    r)   Ú	bessel_j1zPallasKernelOverrides.bessel_j1ü  r*  r+   c                ó   — d| › d| › d�S )Nújax.lax.bessel_i0e(ú) * jnp.exp(jnp.abs(r  r?   rZ   s    r)   Úmodified_bessel_i0z(PallasKernelOverrides.modified_bessel_i0  ó   € ð % Q CÐ';¸A¸3¸bÐAÐAr+   c                ó   — d| › d| › d�S )Nújax.lax.bessel_i1e(r/  r  r?   rZ   s    r)   Úmodified_bessel_i1z(PallasKernelOverrides.modified_bessel_i1  r1  r+   c                ó   — d| › d| › d| › d�S )Nr   z == 0.0, 1.0, jnp.sin(z) / r!   r?   rZ   s    r)   Úspherical_bessel_j0z)PallasKernelOverrides.spherical_bessel_j0  s    € ð ˜A˜3Ð4°Q°C°t¸A¸3¸aÐ@Ð@r+   c                ó   — d| › d�S )Nr.  r!   r?   rZ   s    r)   Úi0ezPallasKernelOverrides.i0e  ó   € ð % Q C qÐ)Ð)r+   c                ó   — d| › d�S )Nr3  r!   r?   rZ   s    r)   Úi1ezPallasKernelOverrides.i1e#  r9  r+   c                ó   — d| › d|› d�S )Nzjax.scipy.special.gammainc(r    r!   r?   ©r[   Úys     r)   ÚgammainczPallasKernelOverrides.gammainc(  s   € ð -¨Q¨C¨r°!°°AÐ6Ð6r+   c                ó   — d| › d|› d�S )Nzjax.scipy.special.gammaincc(r    r!   r?   r=  s     r)   Ú	gammaincczPallasKernelOverrides.gammaincc.  s   € ð .¨a¨S°°1°#°QÐ7Ð7r+   c                ó   — d| › d|› d�S )Nzjax.scipy.special.polygamma(z.astype(jnp.int32), r!   r?   r=  s     r)   Ú	polygammazPallasKernelOverrides.polygamma7  s   € ð .¨a¨SÐ0DÀQÀCÀqÐIÐIr+   c                ó   — d| › d�S )Nzjax.scipy.special.ndtri(r!   r?   rZ   s    r)   ÚndtrizPallasKernelOverrides.ndtri=  s   € ð *¨!¨¨AÐ.Ð.r+   c                ó   — d| › d|› d�S )Nzjax.scipy.special.zeta(r    r!   r?   r=  s     r)   ÚzetazPallasKernelOverrides.zetaB  s   € ð )¨¨¨2¨a¨S°Ð2Ð2r+   c                ó   — d| › d|› d�S )Nzjax.scipy.special.xlogy(r    r!   r?   r=  s     r)   ÚxlogyzPallasKernelOverrides.xlogyG  s   € ð *¨!¨¨B¨q¨c°Ð3Ð3r+   c                ó   — d| › d|› d�S )Nzjax.scipy.special.xlog1py(r    r!   r?   r=  s     r)   Úxlog1pyzPallasKernelOverrides.xlog1pyL  s   € ð ,¨A¨3¨b°°°1Ð5Ð5r+   c                ó>   — d| › d|› d| › d| › d|› d| › d|› d|› d	| › d
�S )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?   ©r[   Úns     r)   Úchebyshev_polynomial_tz,PallasKernelOverrides.chebyshev_polynomial_tQ  s[   € ð !  ð $Ø�cÐ0°°ð 4Ø˜ð Ø�sÐ5°a°Sð 9Ø˜˜M¨!¨Ð,HÈÈÈ:ð	Wð	
r+   c                óÒ   — dj                  g d‘| › ‘d‘|› ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘|› ‘d	‘|› ‘d
‘| › ‘d‘| › ‘d‘| › ‘d‘|› ‘d‘|› ‘d‘|› ‘d‘|› ‘d‘| › ‘d‘| › ‘d‘«      S )NÚ rM  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)))))©ÚjoinrN  s     r)   Úchebyshev_polynomial_uz,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                ó  — dj                  g d‘|› ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘| › ‘d‘|› ‘d	‘| › ‘d
‘| › ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘«      S )NrR  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(ú)))))))rU  rN  s     r)   Úchebyshev_polynomial_vz,PallasKernelOverrides.chebyshev_polynomial_vv  ó
  € ÷)ó )ˆ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                ó  — dj                  g d‘|› ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘| › ‘d	‘|› ‘d
‘| › ‘d‘| › ‘d‘| › ‘d	‘|› ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘«      S )NrR  r   rY  rZ  r[  r`  r]  z**2 + 2*r\  r^  z**3 + 4*r_  ra  z**4 + 8*rb  rc  z	**5 + 16*rd  re  z + 1, jnp.zeros_like(rf  rU  rN  s     r)   Úchebyshev_polynomial_wz,PallasKernelOverrides.chebyshev_polynomial_w‡  rh  r+   c                ó6   — t         j                  d| › d�|«      S ©Nz(2 * z - 1))rX   rP  rN  s     r)   Úshifted_chebyshev_polynomial_tz4PallasKernelOverrides.shifted_chebyshev_polynomial_t˜  ó    € ô %×;Ñ;¸eÀAÀ3ÀeÐ<LÈaÓPÐPr+   c                ó6   — t         j                  d| › d�|«      S rl  )rX   rW  rN  s     r)   Úshifted_chebyshev_polynomial_uz4PallasKernelOverrides.shifted_chebyshev_polynomial_u�  rn  r+   c                ó6   — t         j                  d| › d�|«      S rl  )rX   rg  rN  s     r)   Úshifted_chebyshev_polynomial_vz4PallasKernelOverrides.shifted_chebyshev_polynomial_v¢  rn  r+   c                ó6   — t         j                  d| › d�|«      S rl  )rX   rj  rN  s     r)   Úshifted_chebyshev_polynomial_wz4PallasKernelOverrides.shifted_chebyshev_polynomial_w§  rn  r+   c                óÒ   — dj                  g d‘|› ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘|› ‘d	‘| › ‘d
‘| › ‘d‘|› ‘d‘| › ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘«      S )NrR  r   rY  rZ  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(rf  rU  rN  s     r)   Úhermite_polynomial_hz*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                óÒ   — dj                  g d‘|› ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘|› ‘d	‘| › ‘d
‘| › ‘d‘|› ‘d‘| › ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘«      S )NrR  r   rY  rZ  rS  rv  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 * rw  rf  rU  rN  s     r)   Úhermite_polynomial_hez+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                ó  — dj                  g d‘|› ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘| › ‘d	‘|› ‘d
‘| › ‘d‘| › ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘«      S )NrR  r   rY  rZ  z == 1, 1 - rv  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(rf  rU  rN  s     r)   Úlaguerre_polynomial_lz+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                óÒ   — dj                  g d‘|› ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘|› ‘d	‘| › ‘d
‘| › ‘d‘|› ‘d‘| › ‘d‘| › ‘d‘|› ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘| › ‘d‘«      S )NrR  r   rY  rZ  rS  rv  z == 2, (3 * z**2 - 1) / 2, jnp.where(z == 3, (5 * rz  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(rf  rU  rN  s     r)   Úlegendre_polynomial_pz+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                ó   — d| › d�S )Nzjnp.reciprocal(r!   r?   rZ   s    r)   Ú
reciprocalz PallasKernelOverrides.reciprocalí  r“   r+   c                ó   — d| › d�S )Nzjnp.square(r!   r?   rZ   s    r)   ÚsquarezPallasKernelOverrides.squareñ  s   € à˜Q˜C˜qÐ!Ð!r+   c                ó   — d| › d|› d|› d�S )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((rT  z) + (r  r?   )r˜   r™   r&   s      r)   ÚfmazPallasKernelOverrides.fmaö  s   € ð �A�3�e˜A˜3˜e A 3 bÐ)Ð)r+   c                ó   — d| › d|› d�S )Nzjnp.copysign(r    r!   r?   r—   s     r)   ÚcopysignzPallasKernelOverrides.copysignÿ  s   € ð ˜q˜c  A 3 aÐ(Ð(r+   c                ó   — d| › d|› d�S )Nzjnp.nextafter(r    r!   r?   r—   s     r)   Ú	nextafterzPallasKernelOverrides.nextafter  s   € ð   ˜s " Q C qÐ)Ð)r+   c                ó   — d| › d|› d�S )Nz
jnp.ldexp(r    r!   r?   r—   s     r)   ÚldexpzPallasKernelOverrides.ldexp	  r  r+   c                ó   — d| › d�S )Nz
jnp.frexp(r!   r?   rZ   s    r)   ÚfrexpzPallasKernelOverrides.frexp  rŠ   r+   c                ó   — d| › d�S )Nz	jnp.modf(r!   r?   rZ   s    r)   ÚmodfzPallasKernelOverrides.modf  rt   r+   c                ó   — d| › d|› d�S )Nzjnp.bitwise_and(r    r!   r?   r—   s     r)   Úbitwise_andz!PallasKernelOverrides.bitwise_and  ó   € à! !  B q c¨Ð+Ð+r+   c                ó   — d| › d|› d�S )Nzjnp.bitwise_or(r    r!   r?   r—   s     r)   Ú
bitwise_orz PallasKernelOverrides.bitwise_or  ó   € à    2 a S¨Ð*Ð*r+   c                ó   — d| › d|› d�S )Nzjnp.bitwise_xor(r    r!   r?   r—   s     r)   Úbitwise_xorz!PallasKernelOverrides.bitwise_xor   r“  r+   c                ó   — d| › d�S )Nzjnp.bitwise_not(r!   r?   rZ   s    r)   Úbitwise_notz!PallasKernelOverrides.bitwise_not$  r  r+   c                ó   — d| › d|› d�S )Nzjnp.left_shift(r    r!   r?   r—   s     r)   Ú
left_shiftz PallasKernelOverrides.left_shift(  r–  r+   c                ó   — d| › d|› d�S )Nzjnp.right_shift(r    r!   r?   r—   s     r)   Úright_shiftz!PallasKernelOverrides.right_shift,  r“  r+   c                ó¶   — t         j                  j                  j                  d|«      }dt         j                  j                  j	                  | «      › d|› d�S )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_seedzPallasKernelOverrides.load_seed1  sH   € ô —h‘h—m‘m×/Ñ/Ð0BÀFÓKˆØ”1—8‘8—=‘=×&Ñ& tÓ,Ð-¨V°K°=ÀÐBÐBr+   c                ó   — d| › d|› d|› d�S )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)   ÚrandzPallasKernelOverrides.rand8  s,   € ð@Ø@D¸vð FØ"˜8Ð#DÀVÀHÈHðVð	
r+   c                ó   — d| › d|› d|› d�S )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)   ÚrandnzPallasKernelOverrides.randnH  s,   € ð@Ø@D¸vð FØ"˜8Ð#DÀVÀHÈHðVð	
r+   c                ó&   — d| › d|› d|› d|› d|› d�S )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)   Ú	randint64zPallasKernelOverrides.randint64W  s>   € ð
@Ø@D¸vÐE]Ð^aÐ]bÐbdÐeiÐdjð kØ"˜8Ð#DÀVÀHÈHðVð	
r+   )r[   r:   r9   r:   )r˜   r:   r™   r:   r9   r:   )r¡   r:   r˜   r:   r™   r:   r9   r:   )r­   r:   r®   zCallable[[], str]r¯   r¨   r9   r:   )NT)
r[   r:   r»   útorch.dtyper¼   ztorch.dtype | Noner½   rÛ   r9   r:   )r[   r:   r»   r³  r¼   r³  r9   r:   )r%   r8   r»   r³  r9   r:   )r»   r³  r9   r:   )r[   r:   r  r:   r  r:   r9   r:   )r[   r:   r>  r:   r9   r:   )r[   r:   rO  r:   r9   r:   )r˜   r:   r™   r:   r&   r:   r9   r:   )r£  r:   r¤  r:   r9   r:   )r«  r:   r¤  r:   r9   r:   )
r«  r:   r¤  r:   r°  r:   r±  r:   r9   r:   ){r;   r<   r=   r>   Ústaticmethodr\   r_   ra   rc   rf   rh   rj   rm   ro   rq   rs   rv   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  r
  r  r  r  r  r  Úclipr  r  r  r!  r#  r%  r'  r)  r,  r0  r4  r6  Úi0r8  Úi1r;  r?  rA  ÚigammaÚigammacrC  rE  rG  rI  rK  rP  rW  rg  rj  rm  rp  rr  rt  rx  r|  r~  r€  r‚  r„  r†  rˆ  rŠ  rŒ  rŽ  r�  r’  r•  r˜  rš  rœ  rž  r¥  r¬  r®  r²  r?   r+   r)   rX   rX   n   sF	  „ ñð òó ðð òó ðð òó ðð ò ó ð ð ò ó ð ð ò ó ð ð ò"ó ð"ð ò"ó ð"ð ò"ó ð"ð òó ðð ò ó ð ð ò!ó ð!ð òó ðð ò!ó ð!ð ò ó ð ð ò!ó ð!ð ò ó ð ð ò%ó ð%ð òó ðð òó ðð ò!ó ð!ð ò ó ð ð ò!ó ð!ð ò!ó ð!ð ò&ó ð&ð ò&ó ð&ð ò&ó ð&ð ò(ó ð(ð ò(ó ð(ð ò.ó ð.ð ò;ó ð;ð& ð )-Ø"&ð	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+   rX   c                  óN   — e Zd ZU dZded<   ded<   ded<   ded<   ded	<   ded
<   y)Ú_IndirectAccessInfoz=Describes a detected indirect (data-dependent) buffer access.r:   Útable_paramÚtable_buf_nameÚtupleÚtable_shapeÚintÚindirect_dimÚindirect_varÚindices_paramN©r;   r<   r=   r>   Ú__annotations__r?   r+   r)   r»  r»  b  s)   … áGàÓØÓØÓØÓØÓØÔr+   r»  c                  ó&   — e Zd ZU dZded<   ded<   y)Ú_BufferIndexingzHEncapsulates index string and flattening requirements for buffer access.r:   Ú	index_strrÛ   Úneeds_flattenNrÄ  r?   r+   r)   rÇ  rÇ  n  s   … áRàƒNØÔr+   rÇ  c                  ó:   — e Zd ZU dZded<   ded<   ded<   ded	<   y
)Ú_BroadcastedIterVarzFEncapsulates information needed to codegen a broadcasted iteration varrÀ  Úidxúsympy.SymbolÚvar_symr   Úentryú
int | NoneÚ
length_valNrÄ  r?   r+   r)   rË  rË  v  s   … áPð 
ƒHØÓØÓØÔr+   rË  c                  óÆ   — e Zd ZU dZded<   ded<   ded<   ded<   ded	<   d
ed<   d
ed<   d
ed<   d
ed<   ded<   ded<   d
ed<   d
ed<   d
ed<   d
ed<   ded<   ded<   ded<   y)Ú_CodegenContextz@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_lookupúdict[str, bool]Úaliasable_flagsÚalias_paramsÚpointer_tailÚkernel_input_paramsÚfull_kernel_paramszOrderedSet[str]Únon_alias_out_setú	list[int]Úcopy_output_indicesúlist[tuple[int, int]]Úalias_pairsNrÄ  r?   r+   r)   rÓ  rÓ  �  sp   … áJà
ÓØÓØƒLØÓØÓØÓØÓØÓØÓØ(Ó(Ø$Ó$ØÓØÓØ"Ó"Ø!Ó!Ø&Ó&Ø"Ó"Ø&Ô&r+   rÓ  c                  óŽ  ‡ — e Zd ZU dZeZeZded<   ˆ fd„Z		 	 	 	 	 	 	 	 	 	 dQd„Z
dRd„ZdRd„ZdRd„ZdSd	„ZdTd
„ZdUd„ZdVd„ZdUd„Z	 	 	 	 	 	 dWd„ZedXd„«       Z	 	 	 	 	 	 dYd„ZedZd„«       Zd[d„Z	 	 	 	 	 	 d\d„Z	 	 	 	 	 	 d]d„Zed^d„«       Zd_d„Zed`d„«       Zedad„«       Zedbd„«       Zdcd„Z ddd„Z!ded„Z"dfd„Z#dgd„Z$dhd„Z%	 	 	 	 did „Z&	 	 	 	 	 	 djd!„Z'	 	 	 	 	 	 	 	 	 	 dkd"„Z(	 	 	 	 	 	 	 	 dld#„Z)	 	 	 	 	 	 	 	 dld$„Z*	 	 	 	 	 	 	 	 dld%„Z+edmd&„«       Z,dnd'„Z-	 	 	 	 	 	 	 	 dod(„Z.dpd)„Z/	 	 	 	 	 	 	 	 	 	 dqd*„Z0drd+„Z1	 	 	 	 	 	 	 	 dsd,„Z2	 	 	 	 	 	 dtd-„Z3	 	 	 	 	 	 	 	 dud.„Z4dvd/„Z5	 	 	 	 	 	 	 	 dwd0„Z6	 dx	 	 	 	 	 	 	 	 	 	 	 	 	 dyd1„Z7	 	 	 	 	 	 	 	 	 	 	 	 dzd2„Z8e9jt                  d{d3„«       Z;dRd4„Z<e9jt                  	 dx	 	 	 	 	 	 	 	 	 d|d5„«       Z=e	 d}	 	 	 	 	 	 	 d~d6„«       Z>	 d	 	 	 	 	 d€d7„Z?	 	 	 	 	 	 	 	 d�d8„Z@	 	 	 	 	 	 	 	 d‚d9„ZA	 	 	 	 	 	 	 	 	 	 dƒd:„ZBed„d;„«       ZCdgd<„ZDdxd…d=„ZEe	 	 	 	 	 	 d†d>„«       ZF	 	 	 	 	 	 	 	 d‡d?„ZG	 	 	 	 	 	 	 	 dˆd@„ZHd‰dA„ZIdŠdB„ZJ	 	 d‹dC„ZK	 	 	 	 	 	 dŒdD„ZL	 	 	 	 	 	 d�dE„ZMe	 	 	 	 	 	 	 	 dŽdF„«       ZNd�dG„ZO	 	 	 	 	 	 d�dH„ZPd�dI„ZQd‰dJ„ZR	 	 	 	 	 	 	 	 	 	 d‘dK„ZSd’dL„ZT	 	 	 	 	 	 d’dM„ZU	 	 	 	 	 	 d’dN„ZVe	 	 	 	 	 	 	 	 	 	 d“dO„«       ZWdxd”dP„ZXˆ xZYS )•ÚPallasKernelz½Pallas kernel codegen for TPU and GPU (Mosaic backend).

    Generates Python code that defines a Pallas kernel and a host entrypoint,
    compiled and loaded via async_compile.pallas.
    zCallable[[sympy.Expr], str]rÍ   c                óØ  •— t        ‰| �  |i |¤Ž t        j                  j	                  «       }|j
                  dk(  | _        |j
                  dk(  | _        | j                  | _        g | _	        i | _
        t        «       | _        i | _        i | _        d | _        d | _        g | _        | j$                  j'                  «       D ]B  }|j(                  j*                  D ]'  }| j"                  j-                  |j.                  «       Œ) ŒD t        «       | _        t        «       | _        d| _        i | _        t        «       | _        d | _        i | _        i | _        |j
                  dk(  | _        y )NÚcudar³   F) ÚsuperrN   r   r·   r¸   r¹   Úis_gpurÖ  Úuse_emit_pipelineÚstore_with_outputÚload_index_exprsr   Úoutputs_need_readÚpermuted_input_buffersÚcollapsed_reshape_inputsÚcollapsed_output_shapeÚ_cpu_max_grid_productÚ_output_buffer_namesÚfeaturesÚscheduler_nodesÚread_writesÚwritesÚappendr£  rÈ   Útile_relative_iter_varsÚhas_flatten_indexingÚstrided_input_buffersÚflatten_indexed_buffersÚindirect_accessÚ_cse_to_paramÚ_param_to_graph_name)r$   r#   rQ   ÚdeviceÚsnodeÚdepÚ	__class__s         €r)   rN   zPallasKernel.__init__£  sL  ø€ Ü‰Ñ˜$Ð) &Ò)ä—‘×4Ñ4Ó6ˆØ—k‘k VÑ+ˆŒØ—k‘k UÑ*ˆŒð "&§¡ˆÔà8:ˆÔà79ˆÔä2<³,ˆÔàBDˆÔ#ØDFˆÔ%Ø>BˆÔ#Ø15ˆÔ"ð 02ˆÔ!Ø—]‘]×2Ñ2Ó4ò 	;ˆEØ×(Ñ(×/Ñ/ò ;�Ø×)Ñ)×0Ñ0°·±Õ:ñ;ð	;ô 9C»ˆÔô BLÃˆÔ$à$)ˆÔ!ð
 MOˆÔ"ô 9C»ˆÔ$à;?ˆÔØ-/ˆÔØ46ˆÔ!Ø—k‘k UÑ*ˆ�r+   c                 ó   — y)z)Check array bounds for indirect indexing.Nr?   )r$   r%   ÚsizeÚlowerÚuppers        r)   Úcheck_boundszPallasKernel.check_boundsÓ  ó   � r+   c                óª   — | j                  |«      }t        |t        j                  «      ry|j                  rt        |«      S | j                  |«      S )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_strzPallasKernel._get_index_strÚ  sN   € ð" ×.Ñ.¨uÓ5ˆô �n¤e§l¡lÔ3ØØ×&Ò&ä�~Ó&Ð&ð ×-Ñ-¨nÓ=Ð=r+   c                óì  — | j                   sy| j                  |«      }|j                  t        «      r;| j                  j                  | j                  |«      «       | j                  |«      S t        j                  j                  j                  |«      }| j                  |«      }| j                  j                  |«       t        |«      dk(  rt        |«      S t        |«      dk(  �rt        t        |«      «      }t!        j"                  ||«      }t!        j$                  ||«      }|�—||z
  }t        j                  j                  j                  |«      }|dk  r| j                  |«      S |dk(  ry|dk7  r| j                  |«      S 	 t'        |«      }|dk  r| j                  |«      S 	 | j                  |«      › d�S ||z
  }t        j                  j                  j                  |«      }|dk(  r||k(  ryyt        |«      dkD  rBd}|D ]7  }t!        j"                  ||«      }t!        j$                  ||«      }|dk7  sŒ5d} n |ryyy# t(        t*        f$ r | j                  |«      cY S w xY w)aD  
        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   r·   ÚsizevarsÚsimplifyÚlenr:   ÚnextÚiterr   Úget_subexpr_involving_symbolÚmatch_affine_block_exprrÀ  Ú	TypeErrorÚ
ValueError)	r$   r  Ú	used_varsrÔ   Úvar_exprÚstrider¤  Ú
offset_valÚall_unit_strides	            r)   r  z"PallasKernel._convert_to_jax_sliceý  si  € ð ×ÒØð ×$Ñ$ UÓ+ˆð �9‰9”_Ô%à×Ñ×&Ñ& t×'?Ñ'?ÀÓ'FÔGð —:‘:˜eÓ$Ð$ô —‘× Ñ ×)Ñ)¨%Ó0ˆà×,Ñ,¨UÓ3ˆ	ð 	×Ñ×"Ñ" 9Ô-äˆy‹>˜QÒä�u“:ÐÜ�‹^˜qÓ ä”t˜I“Ó'ˆCô +×GÑGÈÈsÓSˆHô )×@Ñ@ÀÈ3Ó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¨s¢?à ð, ô+ �‹^˜aÒð #ˆOØ ò �Ü.×KÑKÈEÐSVÓW�Ü,×DÑDÀXÈsÓS�Ø˜Q“;Ø&+�OÙðñ àð ð øôA "¤:Ð.ò -ØŸ:™: eÓ,Ò,ð-ús   Æ I É I3É2I3c                óÆ   — |j                   }| j                  «       }||z  }||k7  rt        d|› �«      ‚| j                  j	                  |«       | j                  |«      }|S )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_varsrV   rÈ   rÉ   rÍ   )r$   r  r*  Ú	iter_varsr$  rÈ  s         r)   Ú_generate_strided_indexz$PallasKernel._generate_strided_indexa  sv   € ð ×)Ñ)ˆØ×'Ñ'Ó)ˆ	ð ! 9Ñ,ˆ	Ø˜Ò$ÜØKÈEÈ7ÐSóð ð
 	×Ñ×"Ñ" 9Ô-ð —J‘J˜uÓ%ˆ	ð Ðr+   c                óH   — t        | j                  j                  «       «      S )z*Get the set of iteration variable symbols.)r   Úrange_tree_nodesÚkeys)r$   s    r)   r+  zPallasKernel._get_iter_vars  s   € ä˜$×/Ñ/×4Ñ4Ó6Ó7Ð7r+   c                ó<   — |j                   | j                  «       z  S )z4Get iteration variables used in an index expression.)r*  r+  ©r$   r  s     r)   rÊ   z PallasKernel._get_used_iter_varsƒ  s   € à×!Ñ! D×$7Ñ$7Ó$9Ñ9Ð9r+   c                ó6   — t        | j                  |«      «      S )z7Check if index expression contains iteration variables.)rÛ   rÊ   r2  s     r)   Ú_has_iteration_varsz PallasKernel._has_iteration_vars‡  s   € ä�D×,Ñ,¨UÓ3Ó4Ð4r+   c                óv   — |j                   D �cg c]  }t        |«      j                  d«      sŒ|‘Œ! c}S c c}w )zDGet list of indirect variable symbols (tmp*) in an index expression.Útmp)r*  r:   Ú
startswith)r$   r  Úss      r)   Ú_get_indirect_varszPallasKernel._get_indirect_vars‹  s-   € à ×-Ñ-ÖJ�a´°Q³×1BÑ1BÀ5Õ1I’ÒJÐJùÒJs   �6¯6c                ó<   — t        | j                  |«      «      dkD  S )z6Check if index expression contains indirect variables.r   )r  r9  r2  s     r)   Ú_has_indirect_varszPallasKernel._has_indirect_vars�  s   € ä�4×*Ñ*¨5Ó1Ó2°QÑ6Ð6r+   c                ó|  — | j                  |«      s|j                  t        «      ry|| j                  v ry| j	                  |«      }|€y|\  }}}}}|D �cg c]  }| j                  |«      ‘Œ }}t        d„ |D «       «      ryt        t        t           |«      }t        |«      }	|	dk(  ry| j                  |«      }
| j                  |«      }|syt        |	«      D �cg c]  }ddg‘Œ }}i }t        j                  j                   j#                  |«      }|D ]Û  }t%        j&                  ||«      }t%        j(                  ||«      }|€ y| j                  |«      }|�|dk  r yd}t        |	«      D ](  }|
|   dk(  rŒ||
|   z  dk(  sŒ||
|   z  }|dk\  sŒ&|} n |€ y||
|   z  }|dk  s||   |z  dk7  r y|||   d<   |||<   t        j                  j                   j#                  ||z
  «      }ŒÝ | j                  |«      }|€y|dk  ryt        |	«      D ]!  }|
|   dkD  sŒ||
|   z  ||   d<   ||
|   z  }Œ# |dk7  ryt+        d„ |D «       «      ryg }d}t        |	«      D ]?  }||   \  }}||z  }||z  }||   |z  }||k\  r y|||z
  z  }|j-                  |||f«       ŒA | j/                  |«      \  }}||k7  ry|j1                  «       D ]\  \  }}|| j2                  vr y| j                  | j2                  |   j4                  «      }|€ y||   \  } }!}"||   | z  |"z
  }#||#k7  sŒ\ y |S c c}w c c}w )a'  Decompose a flat index into per-dimension (stride, offset, skip) triples.

        Given flat index like ``64*x0 + 2*x1 + 5`` and buffer shape ``(32, 64)``
        with C-contiguous strides ``[64, 1]``:
          - x0 coefficient 64 / buffer_stride[0]=64 -> dim 0: stride=1
          - x1 coefficient 2 / buffer_stride[1]=1  -> dim 1: stride=2
          - constant 5: dim 0 gets 5//64=0, dim 1 gets 5//1=5
          - dim 1 offset 5 with stride 2: skip=5//2=2, offset=5%2=1

        Returns per-dim ``[(stride, offset, skip), ...]`` where:
          - stride: access stride on this dim (1 = contiguous)
          - offset: static index into the stride dim (0 <= offset < stride)
          - skip: number of stride-blocks to skip at the start of this dim
        Returns None if decomposition fails.
        Nc              3  ó2   K  — | ]  }|d u xs |dk  –— Œ y ­w©Nr   r?   ©Ú.0r8  s     r)   ú	<genexpr>z9PallasKernel._decompose_strided_access.<locals>.<genexpr>³  s    è ø€ Ò> qˆq�DˆyÒ"˜A ™FÓ"Ñ>ùó   ‚r   r   c              3  ó,   K  — | ]  \  }}|d k(  –— Œ y­w©r   Nr?   )r@  r8  Ú_s      r)   rA  z9PallasKernel._decompose_strided_access.<locals>.<genexpr>ö  s   è ø€ Ò)™$˜!˜Qˆq�A�vÑ)ùs   ‚)r;  r  r
   r  Ú_get_buffer_infoÚ	_safe_intÚanyr   ÚlistrÀ  r  Ú_c_contiguous_stridesrÊ   Úranger   r·   r  r  r   r   r!  Úallrý  Ú _compute_output_numel_from_indexÚitemsr/  Úlength)$r$   r  r£  rM   rE  Úbuf_sizer8  Úbuf_shape_or_noneÚ	buf_shapeÚndimÚ	c_stridesr$  r1   Ú
var_to_dimÚ	remainingrÔ   r%  ÚcoeffÚ	coeff_intÚdimÚdÚper_dim_strider'  Ú
decomposedÚoutput_numel_expectedr&  Ú
raw_offsetr¤  ÚskipÚn_blocksÚoutput_numelÚ	var_rangeÚstride_dÚ	_offset_dÚskip_dÚeffective_sizes$                                       r)   Ú_decompose_strided_accessz&PallasKernel._decompose_strided_access“  sî  € ð$ ×"Ñ" 5Ô)¨U¯Y©Y´Ô-GØð �4×/Ñ/Ñ/Øà×$Ñ$ TÓ*ˆØˆ<ØØ#Ñˆˆ8�Q˜˜1à8@ÖA°1˜TŸ^™^¨AÕ.ÐAÐÐAÜÑ>Ð,=Ô>Ô>ØÜ#¤D¬¡IÐ/@ÓAˆ	Ü�9‹~ˆØ�1Š9Øà×.Ñ.¨yÓ9ˆ	ð ×,Ñ,¨UÓ3ˆ	ÙØô 49¸³;Ö"?¨a A q¢6Ð"?ˆÐ"?à.0ˆ
ä—G‘G×$Ñ$×-Ñ-¨eÓ4ˆ	Øò 	HˆCÜ*×GÑGÈ	ÐSVÓWˆHÜ'×?Ñ?ÀÈ#ÓNˆEØˆ}ÙØŸ™ uÓ-ˆIØÐ  I°¢NÙð ˆCÜ˜4“[ò �Ø˜Q‘< 1Ò$ØØ˜y¨™|Ñ+¨qÓ0Ø%.°)¸A±,Ñ%>�NØ%¨Ó*Ø˜Ùðð ˆ{Ùà&¨)°C©.Ñ8ˆNØ Ò! Y¨s¡^°nÑ%DÈÒ%IÙà+ˆF�3‰K˜‰NØ!ˆJ�s‰OÜŸ™×(Ñ(×1Ñ1°)¸hÑ2FÓG‰Ið9	Hð@ —^‘^ IÓ.ˆ
ØÐØØ˜Š>ØÜ�t“ò 	7ˆAØ˜‰|˜aÓØ)¨Y°q©\Ñ9��q‘	˜!‘Ø'¨)°A©,Ñ6‘
ð	7ð ˜Š?Øô
 Ñ) &Ô)Ô)Øð
 24ˆ
Ø !ÐÜ�t“ò 	6ˆAØ!'¨¡ÑˆF�JØ &Ñ(ˆFØ Ñ'ˆDØ  ‘| vÑ-ˆHØ�xÒÙØ! X°¡_Ñ4Ð!Ø×Ñ˜v v¨tÐ4Õ5ð	6ð ×?Ñ?ÀÓF‰ˆ�aØÐ0Ò0Øð #×(Ñ(Ó*ò 		‰HˆC�Ø˜$×/Ñ/Ñ/ÙØŸ™ t×'<Ñ'<¸SÑ'A×'HÑ'HÓIˆIØÐ ÙØ*4°S©/Ñ'ˆH�i Ø& s™^¨xÑ7¸&Ñ@ˆNØ˜NÓ*Ùð		ð ÐùòU Bùò  #@s   ÁL4ÃL9c                óÎ   — g }|D ]H  \  }}}|dk(  r|j                  d«       Œ|j                  d«       |j                  t        |«      «       ŒJ | › ddj                  |«      › d�S )z?Build ``buf[:, :, offset]`` for strided dims, ``:`` for others.r   ú:ú[r    ú])rý  r:   rV  )ÚbufÚdecompÚpartsr&  r¤  Ú_skips         r)   Ú_strided_load_exprzPallasKernel._strided_load_expr  sm   € ð ˆØ%+ò 	*Ñ!ˆF�F˜EØ˜Š{Ø—‘˜SÕ!à—‘˜SÔ!Ø—‘œS ›[Õ)ð	*ð ��a˜Ÿ	™	 %Ó(Ð)¨Ð+Ð+r+   c           
     ó,  — |D �]Ž  }| j                  |«      }|�|| j                  vrŒ&| j                  |   }| j                  |«      }|€ŒI|\  }}}}}g }	t        |«      D ]x  \  }
\  }}}| j	                  ||
   «      }|€ Œ€|dkD  r8|	j                  t        ||z  «      «       |	j                  t        |«      «       Œ_|	j                  t        |«      «       Œz |j                  |› d|› ddj                  |	«      › d�«       t        d„ |D «       «      s�Œg }|D ]I  \  }}}|dk(  r|j                  d«       Œ|j                  |d	kD  r|› d�nd«       |j                  d«       ŒK |j                  |› d|› d
dj                  |«      › d�«       �Œ‘ y)a	  Emit reshape + optional slice for strided input parameters.

        For each strided param, reshapes ``(M, N)`` to ``(M, N/stride, stride)``
        and, when ``skip > 0``, slices off leading blocks so the remaining
        elements align with the output.
        Nr   ú = ú	.reshape(r    r!   c              3  ó.   K  — | ]  \  }}}|d kD  –— Œ y­w©r   Nr?   )r@  rE  r_  s      r)   rA  z9PallasKernel._codegen_strided_reshapes.<locals>.<genexpr>J  s   è ø€ Ò:¡J A q¨$�t˜a•xÑ:ùs   ‚ri  r   rj  rk  )
Ú_param_to_buf_namer   rF  Ú	enumeraterG  rý  r:   Ú	writelinerV  rH  )r$   rÔ  ÚparamsÚparamÚbuf_nameÚstridesrM   rE  rP  Únew_shape_partsrZ  r&  Ú_offsetro  rY  Úslice_partsr_  s                    r)   Ú_codegen_strided_reshapesz&PallasKernel._codegen_strided_reshapes*  s¹  € ð ó 	SˆEØ×.Ñ.¨uÓ5ˆHØÐ 8°4×3MÑ3MÑ#MØØ×0Ñ0°Ñ:ˆGØ×(Ñ(¨Ó2ˆDØˆ|ØØ#'Ñ ˆAˆx˜˜A˜qØ)+ˆOÜ/8¸Ó/Aò SÑ+�Ñ+�F˜G UØ—n‘n X¨a¡[Ó1�Ø�;ÙØ˜A’:Ø#×*Ñ*¬3¨s°f©}Ó+=Ô>Ø#×*Ñ*¬3¨v«;Õ7à#×*Ñ*¬3¨s«8Õ4ðSð —‘Ø�g˜S   y°·±¸?Ó1KÐ0LÈAÐNôô Ñ:°'Ô:Ö:Ø-/�KØ18ò 4Ñ-˜ ¨Ø! Qš;Ø'×.Ñ.¨sÕ3à'×.Ñ.¸TÀAºX°$°°q©zÈ3ÔOØ'×.Ñ.¨sÕ3ð4ð —N‘N e W¨C°¨w°a¸¿	¹	À+Ó8NÐ7OÈqÐ#QÖRñ?	Sr+   c                ó  —  t        | dd„ «      «       }|€yt        |dd«      }|�t        |«      |k7  ryg }|D ]D  }t        |t        t        j
                  f«      rt        |«      nd}|€ y|j                  |«       ŒF |S )z5Extract actual output buffer strides from its layout.Ú
get_layoutc                  ó   — y rS   r?   r?   r+   r)   ú<lambda>z6PallasKernel._get_actual_out_strides.<locals>.<lambda>W  r  r+   Nr&  )Úgetattrr  r§   rÀ  r  ÚIntegerrý  )Úout_bufrO  ÚlayoutÚ
stride_rawr|  r8  Úvs          r)   Ú_get_actual_out_stridesz$PallasKernel._get_actual_out_stridesT  s�   € ð >”˜ ,±Ó=Ó?ˆØˆ>ØÜ˜V X¨tÓ4ˆ
ØÐ¤ Z£°AÒ!5ØØˆØò 	ˆAÜ$ Q¬¬e¯m©mÐ(<Ô=”�A”À4ˆAØˆyÙØ�N‰N˜1Õð		ð
 ˆr+   c                ó|  — i }g }|D ]K  }| j                   |   }|j                  }||vrg ||<   |j                  |«       ||   j                  |«       ŒM g }t        |«      D ]  }|j	                  ||   «       Œ i }d}	|D ]8  }| j                  | j                   |   j                  «      }
|
€ y|	||<   |	|
z  }	Œ: |S )a  Compute store-side linearization coefficients from range tree nesting.

        The tree structure encodes the output iteration order: later
        trees (prefix ``x``) are innermost, earlier trees (``y``, ``z``)
        are outer.  Within a tree, dict order goes inner-to-outer.
        The innermost variable gets coefficient 1; each successive
        variable (moving outward) multiplies by the previous range.

        Returns ``{sympy.Symbol: int}`` mapping each RT var to its store
        coefficient, or ``None`` on failure.
        r   N)r/  Úprefixrý  ÚreversedÚextendrG  rO  )r$   ÚorderedÚprefix_groupsÚprefix_orderrŠ  Únoder'   Úinner_to_outerÚcoeffsrW  Úszs              r)   Ú_compute_store_coeffsz"PallasKernel._compute_store_coeffse  sð   € ð *,ˆØ"$ˆØò 	'ˆAØ×(Ñ(¨Ñ+ˆDØ—‘ˆAØ˜Ñ%Ø#%�˜aÑ Ø×#Ñ# AÔ&Ø˜!Ñ×#Ñ# AÕ&ð	'ð  "ˆÜ˜,Ó'ò 	4ˆAØ×!Ñ! -°Ñ"2Õ3ð	4àˆØˆØò 	ˆAØ—‘ × 5Ñ 5°aÑ 8× ?Ñ ?Ó@ˆBØˆzÙØˆF�1‰IØ�R‰K‰Eð	ð ˆr+   c                óø  — | j                  |«      }|sy|\  }}}}}|D �cg c]  }| j                  |«      ‘Œ }}t        |«      dk  sd|v ryt        t        t
           |«      }	|sy| j                  |«      }
| j                  j                  «       D ��cg c]  \  }}||
v r|j                  s|‘Œ }}}t        |«      t        |	«      k7  �rGt        |	«      }t        |«      dk(  �r,| j                  | j                  |d      j                  «      t        j                  |	«      k(  rë| j                  |	«      }| j                  D ]Ë  }t        j                   j#                  |«      }|€Œ%|j%                  «       D �cg c]  }| j                  |«      ‘Œ }}t'        d„ |D «       «      st        |«      |k7  rŒr| j)                  ||«      }|€ y| j+                  ||«      }|€ yt	        |«      t	        t-        |«      «      k(  r yt/        |«      c S  y|D �cg c];  }| j1                  t        j                   j2                  j5                  |«      |«      ‘Œ= }}t7        d„ |D «       «      syt        t        t
           |«      }t        |«      }| j                  |	«      }| j9                  |«      }|��(| j                  D �]  }t        j                   j#                  |«      }|€Œ&|j%                  «       D �cg c]  }| j                  |«      ‘Œ }}t'        d„ |D «       «      st        |«      |k7  rŒs| j)                  ||«      }|�’| j+                  |D �cg c]  }||   ‘Œ	 c}|«      }| j+                  t	        |«      |«      }|�R|�Pdg|z  }t-        |«      D ]  }||   |||   <   Œ t	        |«      t	        t-        |«      «      k(  r yt/        |«      c S  y yc c}w c c}}w c c}w c c}w c c}w c c}w )a  Return permutation for a full-array load, or None.

        Computes the permutation by mapping each range-tree variable to
        both an output dimension (via store coefficients + actual output
        strides) and an input dimension (via load coefficients + input
        C-contiguous strides).  The permutation is then:

            perm[out_dim] = in_dim   for each RT variable

        Using actual output strides (not C-contiguous) is critical: the
        scheduler may choose a non-standard output layout (e.g. column-
        major) to optimise for transposed inputs.

        When all dimensions collapse to a single flat RT variable (e.g.
        (2,2,2,2,2) with all dims size 2), infers the permutation
        directly from output strides vs input C-contiguous strides.
        Nr   r   r   c              3  ó$   K  — | ]  }|d u –— Œ
 y ­wrS   r?   r?  s     r)   rA  z:PallasKernel._get_full_load_permutation.<locals>.<genexpr>¼  s   è ø€ Ò8¨˜1 œ9Ñ8ùó   ‚c              3  óJ   K  — | ]  }t        |t        «      xr |d kD  –— Œ y­wru  ©r§   rÀ  ©r@  r&   s     r)   rA  z:PallasKernel._get_full_load_permutation.<locals>.<genexpr>Î  s#   è ø€ ÒD°A”:˜a¤Ó%Ò/¨!¨a©%Ó/ÑDùó   ‚!#c              3  ó$   K  — | ]  }|d u –— Œ
 y ­wrS   r?   r?  s     r)   rA  z:PallasKernel._get_full_load_permutation.<locals>.<genexpr>Ý  s   è ø€ Ò4 Q�q˜D”yÑ4ùrš  )rF  rG  r  r   rI  rÀ  rÊ   r/  rN  Úis_reductionrO  r©   ÚprodrJ  rø  r   r·   Ú
get_bufferÚget_sizerH  r‹  Ú_map_coeffs_to_dimsrK  r¾  Ú_get_index_coefficientr  r  rL  r—  )r$   r£  r  rM   rE  rP  Úis_contiguousr8  Úin_shape_rawÚin_shapeÚ	iter_usedÚer�  rO  Ú
in_stridesÚout_namer‡  Ú	out_shapeÚactualÚpermrŠ  Ú
coeffs_rawr•  Ústore_coeffsÚ	rt_to_outÚrt_to_inÚks                              r)   Ú_get_full_load_permutationz'PallasKernel._get_full_load_permutation‡  sð  € ð( ×$Ñ$ TÓ*ˆÙØØ+/Ñ(ˆˆ8�Q˜˜=Ø3;Ö<¨a˜Ÿ™ qÕ)Ð<ˆÐ<Üˆ|Ó˜qÒ  D¨LÑ$8ØÜ"¤4¬¡9¨lÓ;ˆÙØð ×,Ñ,¨UÓ3ˆ	ð ×-Ñ-×3Ñ3Ó5÷
á��1Ø�I‰~ a§n¢nò ð
ˆñ 
ô
 ˆw‹<œ3˜x›=Ó(ô
 �H“ˆAÜ�7‹|˜qÓ  T§^¡^Ø×%Ñ% g¨a¡jÑ1×8Ñ8ó&ä—‘˜8Ó$ò&%ð "×7Ñ7¸ÓA�
Ø $× 9Ñ 9ò '�HÜŸg™g×0Ñ0°Ó:�GØ�Ø Ø<C×<LÑ<LÓ<NÖ O°q §¡°Õ!2Ð O�IÐ OÜÑ8¨iÔ8Ô8¼CÀ	»NÈaÒ<OØ Ø!×9Ñ9¸'À1ÓE�FØ�~Øð ð  ×3Ñ3°F¸JÓG�DØ�|Øð ô ˜D“z¤T¬%°«(£^Ò3Ù#Ü  ›;Ò&ð#'ð$ ð ö
àð ×'Ñ'¬¯©×(8Ñ(8×(AÑ(AÀ%Ó(HÈ!ÕLð
ˆ
ð 
ô ÑD¸ÔDÔDØÜ ¤¤c¡¨JÓ7ˆä�‹LˆØ×/Ñ/°Ó9ˆ
Ø×1Ñ1°'Ó:ˆð Ñ#Ø ×5Ñ5ó �ÜŸ'™'×,Ñ,¨XÓ6�Ø�?ØØ8?×8HÑ8HÓ8JÖK°1˜TŸ^™^¨AÕ.ÐK�	ÐKÜÑ4¨)Ô4Ô4¼¸I»È!Ò8KØà×5Ñ5°g¸qÓA�ØÐ%Ø $× 8Ñ 8Ø29Ö:¨Q˜ a›Ò:¸Fó!�Ið  $×7Ñ7¼¸V»ÀjÓQ�HØ Ð,°Ð1EØ !˜s Q™w˜Ü!& q£ò =˜AØ19¸!±˜D ¨1¡Ò.ð=ä ›:¬¬e°A«h«Ò7Ù#'Ü$ T›{Ò*Øàð/ð. ùòa =ùó
ùò& !Pùò
ùò$ Lùò ;s$   ¡OÂO"Å<O(ÈA O-Ë8O2ÍO7
c                óp  — | j                  |«      }|sy|\  }}}}}|D �cg c]  }| j                  |«      ‘Œ }}t        |«      dk  sd|v ryt        t        t
           |«      }	|sy| j                  |«      }
| j                  j                  «       D ��cg c]  \  }}||
v r|j                  s|‘Œ }}}t        |«      }|dk  s|t        |	«      k\  ry|D �cg c]*  }| j                  | j                  |   j                  «      ‘Œ, }}d|v ryt        t        t
           |«      }t        j                  |«      t        j                  |	«      k7  ry| j                  |	|«      }|€ydg|z  }d}t        |dz
  dd«      D ]  }|||<   |||   z  }Œ t        j                   j"                  j%                  |«      }|D �cg c]  }| j'                  ||«      ‘Œ }}t)        d„ |D «       «      syt        t        t
           |«      }t+        |«      D ��ci c]  \  }}||“Œ
 }}}g }|D ](  }|j-                  |«      }|€ y|j/                  |«       Œ* | j1                  |«      }|€y| j2                  D �]ž  }t        j                   j5                  |«      }|€Œ&|j7                  «       D �cg c]  }| j                  |«      ‘Œ }}t9        d„ |D «       «      st        |«      dk  rŒst        t        t
           |«      } t        j                  | «      t        j                  |	«      k7  rŒ¶| j                  | t	        |«      «      }!|!€ŒÔdg|z  }"d}t        |dz
  dd«      D ]  }||"|<   ||!|   z  }Œ t+        |"«      D �#�ci c]  \  }#}||#“Œ
 }$}#}g }%|D ],  }|$j-                  ||   «      }#|#€  y|%j/                  |#«       Œ. dg|z  }&t        |«      D ]  }'||'   |&|%|'   <   Œ |&t	        t        |«      «      k(  r yt;        |«      t;        |&«      fc S  yc c}w c c}}w c c}w c c}w c c}}w c c}w c c}}#w )aë  Handle permutation when range tree has collapsed dimensions.

        When simplify_and_reorder merges contiguous dims, the range tree
        has fewer variables than the buffer's rank.  This method detects
        the permutation in the collapsed space and returns
        (collapsed_input_shape, perm) so the caller can generate:
            jnp.permute_dims(load.reshape(collapsed_shape), perm)

        Uses index coefficients on both sides: load-index coefficients
        map vars to collapsed input dims, and store-side coefficients
        (derived from the range tree nesting) map vars to collapsed
        output dims.  Both sets of strides are always unique, so
        matching is unambiguous even with duplicate group sizes.
        Nr   r   r   éÿÿÿÿc              3  óJ   K  — | ]  }t        |t        «      xr |d kD  –— Œ y­wru  rœ  r�  s     r)   rA  z?PallasKernel._get_collapsed_load_permutation.<locals>.<genexpr>.  s#   è ø€ ÒG°A”:˜a¤Ó%Ò/¨!¨a©%Ó/ÑGùrž  c              3  ó$   K  — | ]  }|d u –— Œ
 y ­wrS   r?   r?  s     r)   rA  z?PallasKernel._get_collapsed_load_permutation.<locals>.<genexpr>D  s   è ø€ Ò4 �1˜”9Ñ4ùrš  )rF  rG  r  r   rI  rÀ  rÊ   r/  rN  r   rO  r©   r¡  Ú_group_dims_to_rangesrK  r   r·   r  r  r¥  rL  rw  Úgetrý  r—  rø  r¢  r£  rH  r¾  )(r$   r£  r  rM   rE  rP  r¦  r8  r§  r¨  r©  rª  r�  rO  rŠ  Ú
ranges_rawÚrangesÚ	in_groupsÚcollapsed_in_stridesr&  ÚiÚ
simplifiedÚin_coeffs_rawÚ	in_coeffsÚin_stride_to_dimÚvar_to_in_dimrW  rY  r±  r¬  r‡  Úout_shape_rawr­  Ú
out_groupsÚcollapsed_out_stridesÚjÚout_stride_to_dimÚvar_to_out_dimr¯  r´  s(                                           r)   Ú_get_collapsed_load_permutationz,PallasKernel._get_collapsed_load_permutationñ  sv  € ð" ×$Ñ$ TÓ*ˆÙØØ+/Ñ(ˆˆ8�Q˜˜=Ø3;Ö<¨a˜Ÿ™ qÕ)Ð<ˆÐ<Üˆ|Ó˜qÒ  D¨LÑ$8ØÜ"¤4¬¡9¨lÓ;ˆÙØà×,Ñ,¨UÓ3ˆ	ð ×-Ñ-×3Ñ3Ó5÷
á��1Ø�I‰~ a§n¢nò ð
ˆñ 
ô
 �‹LˆØˆqŠ5�Aœ˜X›Ò&ØØOVÖWÈ!�d—n‘n T×%:Ñ%:¸1Ñ%=×%DÑ%DÕEÐWˆ
ÐWØ�:ÑØÜ ¤¤c¡¨JÓ7ˆÜ�9‰9�VÓ¤§	¡	¨(Ó 3Ò3Øð ×.Ñ.¨x¸Ó@ˆ	ØÐØð !"˜s Q™wÐØˆÜ�q˜1‘u˜b "Ó%ò 	#ˆAØ&,Ð  Ñ#Ø�i ‘lÑ"‰Fð	#ô —W‘W×%Ñ%×.Ñ.¨uÓ5ˆ
ØMTÖUÈ˜×4Ñ4°ZÀÕCÐUˆÐUÜÑG¸ÔGÔGØÜ#¤D¬¡I¨}Ó=ˆ	ä-6Ð7KÓ-L×M¡T Q¨˜A˜q™DÐMÐÑMØˆØò 	&ˆEØ"×&Ñ& uÓ-ˆCØˆ{ÙØ× Ñ  Õ%ð		&ð ×1Ñ1°'Ó:ˆØÐØð ×1Ñ1ó #	3ˆHÜ—g‘g×(Ñ(¨Ó2ˆGØˆØØ8?×8HÑ8HÓ8JÖK°1˜TŸ^™^¨AÕ.ÐKˆMÐKÜÑ4 mÔ4Ô4¼¸MÓ8JÈQÒ8NØÜ#'¬¬S©	°=Ó#AˆIÜ�y‰y˜Ó#¤t§y¡y°Ó':Ò:ØØ×3Ñ3°I¼tÀI»ÓOˆJØÐ!Øð &' C¨!¡GÐ!ØˆFÜ˜1˜q™5 " bÓ)ò (�Ø+1Ð% aÑ(Ø˜* Q™-Ñ'‘ð(ô 3<Ð<QÓ2R× S©$¨!¨Q  A¡Ð SÐÑ SØˆNØò )�Ø%×)Ñ)¨,°q©/Ó:�Ø�9ÚØ×%Ñ% aÕ(ð	)ð �3˜‘7ˆDÜ˜1“Xò ;�Ø*7¸Ñ*:��^ AÑ&Ò'ð;à”tœE !›H“~Ò%ÙÜ˜)Ó$¤e¨D£kÐ2Ò2ðG#	3ðH ùò{ =ùó
ùò Xùò. Vùó
 Nùò" Lùó" !Ts)   ¡PÂPÃ/PÆ=P"ÈP'Ê'P-Í;P2c                ó*  — t        |«      }g }d}t        t        | «      dz
  dd«      D ]@  }|| |   z  }	 |j                  |«      }|j                  |«       |j                  |«       d}ŒB |dk7  s|ry|j                  «        |S # t        $ r Y Œhw xY w)z•Group consecutive dims (right-to-left) to match range values.

        Returns collapsed shape (left-to-right) or None if no valid grouping.
        r   r·  N)rI  rK  r  r  r#  rý  ÚpopÚreverse)Údimsr½  Ú	availableÚgroupsÚproductrÀ  rÌ  s          r)   rº  z"PallasKernel._group_dims_to_rangese  s¨   € ô ˜“Lˆ	ØˆØˆÜ”s˜4“y 1‘} b¨"Ó-ò 	ˆAØ�t˜A‘wÑˆGðØ—o‘o gÓ.�ð �M‰M˜'Ô"Ø�M‰M˜#ÔØ‰Gð	ð �aŠ<™9ØØ�‰ÔØˆøô ò Ùðús   µBÂ	BÂBc                óŽ  — | j                  |«      }| j                  |«      }|r|rt        | j                  |«      d¬«      S |rt        | j	                  |«      d¬«      S | j                  |«      }|j                  t        «      xr |dk7  }|s*|dk7  r%d|v s!|j                  d«      j                  «       sd}t        ||¬«      S )z@Get the index expression string and whether it needs flattening.T©rÈ  rÉ  Fr  ú::ú-)
r;  r4  rÇ  Ú_handle_mixed_indexingrÍ   r  r  r
   ÚlstripÚisdigit)r$   r  Úhas_indirectÚhas_iter_varsrÈ  rÉ  s         r)   Ú_get_index_exprzPallasKernel._get_index_expr|  sÅ   € à×.Ñ.¨uÓ5ˆØ×0Ñ0°Ó7ˆá™MÜ"Ø×5Ñ5°eÓ<ÈDôð ñ Ü"¨T¯Z©Z¸Ó->ÈeÔTÐTà×+Ñ+¨EÓ2ˆIð "ŸI™I¤oÓ6ÒM¸9ÈÑ;MˆMñ ! Y°%Ò%7à 	Ñ)¨Y×-=Ñ-=¸cÓ-B×-JÑ-JÔ-LØ$(�MÜ"¨YÀmÔTÐTr+   c                óD   — 	 t        | «      S # t        t        f$ r Y yw xY w)z0Convert value to int, returning None on failure.N)rÀ  r"  r#  )rÜ   s    r)   rG  zPallasKernel._safe_int”  s'   € ð	Ü�s“8ˆOøÜœ:Ð&ò 	Ùð	ús   ‚
 �žc                óz   — t        | «      }dg|z  }t        |dz
  dd«      D ]  }||dz      | |dz      z  ||<   Œ |S )z0Return C-contiguous strides for the given shape.r   r   r·  )r  rK  )ÚshaperO  r|  rÀ  s       r)   rJ  z"PallasKernel._c_contiguous_stridesœ  sW   € ô �‹JˆØ�#˜‘'ˆÜ�q˜1‘u˜b "Ó%ò 	7ˆAØ   Q¡™¨%°°A±©,Ñ6ˆG�AŠJð	7àˆr+   c                óæ   — i }t        |«      D ]  \  }}||v r y|||<   Œ g }| D ](  }|j                  |«      }|€ y|j                  |«       Œ* t        t	        |«      «      t        | «      k7  ry|S )zÚMap coefficient values to dimension indices via stride matching.

        Returns a list where entry k is the dimension whose stride equals
        coeffs[k], or None if the mapping is ambiguous or incomplete.
        N)rw  r»  rý  r  r   )r•  r|  Ústride_to_dimrZ  r8  Úmappingr&   s          r)   r¤  z PallasKernel._map_coeffs_to_dims¥  s�   € ð )+ˆÜ˜gÓ&ò 	!‰DˆAˆqØ�MÑ!ÙØ ˆM˜!Òð	!ð  ˆØò 	ˆAØ×!Ñ! !Ó$ˆAØˆyÙØ�N‰N˜1Õð		ô
 Œz˜'Ó"Ó#¤s¨6£{Ò2ØØˆr+   c                ó  — d}|j                   j                  «       D ]^  }t        j                  j	                  |«      }|€d}Œ'|j                  «       D ]%  }| j                  |«      }|dk(  rd|fc c S |�Œ$d}Œ' Œ` d|fS )z9Return whether an output has a zero or unknown dimension.FTr   )rÞ  Úvaluesr   r·   Útry_get_bufferr£  rG  )r$   ÚctxÚhas_unknown_dimr{  rl  rY  Údim_ints          r)   Ú_zero_dim_output_flagsz#PallasKernel._zero_dim_output_flags»  s’   € àˆØ×0Ñ0×7Ñ7Ó9ò 
	+ˆHÜ—'‘'×(Ñ(¨Ó2ˆCØˆ{Ø"&�ØØ—|‘|“~ò +�ØŸ.™.¨Ó-�Ø˜a’<Ø Ð0Ô0Ø‘?Ø&*‘Oñ+ð
	+ð �oÐ%Ð%r+   c                óP  ‡‡‡‡‡— | j                   sy| j                  j                  «       D ��cg c]  \  }}|j                  sŒ|‘Œ }}}| j                  j                  «       D ��cg c]  \  }}|j                  rŒ|‘Œ }}}|r|sy| j	                  «       }|r|dk  ry| j                   j                  «       D �]D  \  }}| j                  |«      }|€Œ|\  }	}
}	}}	t        |
«      Š‰dk  rŒ4|D �cg c]  }| j                  |«      ‘Œ }}t        d„ |D «       «      rŒet        t        t           |«      Š|}|D ]  }|j                  |d«      }Œ t        «       Št        j                  j!                  |«      D ]M  }|j"                  rŒ|j%                  «       \  }}	| j                  |«      }|€Œ7|dkD  sŒ=‰j'                  |«       ŒO ‰s�Œt)        ˆˆfd„t+        ‰«      D «       «      }|s�Œ6t        ‰«      dkD  r=t        ‰«      t        |«      k(  r$t        |«      dkD  rt-        ˆfd„|D «       «      c S �Œ�t/        t1        ‰«      «      Š|dz
  ‰z  Št3        ˆfd„t+        ‰dz
  «      D «       «      }|rF|d	   }|}|dkD  r!‰‰|dz
     kD  r|dz  }|dkD  r‰‰|dz
     kD  rŒt        t+        ||dz   «      «      }n,t)        ˆˆˆfd
„t+        ‰«      D «       «      }|st        |«      }t-        ˆfd„|D «       «      c S  t/        t1        | j                   j5                  «       «      «      }|j7                  |d   «      }|dk7  r| j                  |«      ndŠ‰€dŠ|j7                  |d   «      }|dk7  r| j                  |«      nd}|€d}‰|kD  ryyc c}}w c c}}w c c}w )ah  Determine which axes of the loaded array are reduction axes.

        Finds the innermost reduction stride from the load index
        expression, then walks outward through the buffer's dims
        using stride ratios until the accumulated product reaches
        red_numel.  Falls back to stride-direction analysis for
        gather/flatten loads.
        )r·  r   r   c              3  ó$   K  — | ]  }|d u –— Œ
 y ­wrS   r?   r?  s     r)   rA  z3PallasKernel._get_reduction_axes.<locals>.<genexpr>é  s   è ø€ Ò6 �1˜”9Ñ6ùrš  r   c              3  ó2   •K  — | ]  }‰|   ‰v sŒ|–— Œ y ­wrS   r?   )r@  rÀ  Úr_coeffsr|  s     €€r)   rA  z3PallasKernel._get_reduction_axes.<locals>.<genexpr>þ  s   øè ø€ Ò@�q¨°©°xÒ)?”Ñ@ùs   ƒ�c              3  ó(   •K  — | ]	  }|‰z
  –— Œ y ­wrS   r?   ©r@  rÀ  Únds     €r)   rA  z3PallasKernel._get_reduction_axes.<locals>.<genexpr>  s   øè ø€ Ò 9¨A  R¥Ñ 9ùó   ƒc              3  ó:   •K  — | ]  }‰|   ‰|d z      kD  –— Œ y­wrD  r?   )r@  rÀ  r|  s     €r)   rA  z3PallasKernel._get_reduction_axes.<locals>.<genexpr>  s"   øè ø€ ÒSÀ ¨¡
¨W°Q¸±U©^Õ ;ÑSùó   ƒr·  c              3  óH   •K  — | ]  }‰‰|   k  r‰|   ‰‰z   k  r|–— Œ y ­wrS   r?   )r@  rÀ  Úr_strideÚspanr|  s     €€€r)   rA  z3PallasKernel._get_reduction_axes.<locals>.<genexpr>  s4   øè ø€ ò àØ 7¨1¡:Ò-°'¸!±*¸tÀh¹Ò2Nô ñùó   ƒ"c              3  ó(   •K  — | ]	  }|‰z
  –— Œ y ­wrS   r?   rð  s     €r)   rA  z3PallasKernel._get_reduction_axes.<locals>.<genexpr>!  s   øè ø€ Ò. A˜˜R�Ñ.ùrò  ©r   )rò  r/  rN  r   Ú_compute_reduction_numelrF  r  rG  rH  r   rI  rÀ  Úsubsr   r  ÚAddÚ	make_argsÚ	is_numberÚas_coeff_MulÚaddÚsortedrK  r¾  r  r  rL  rå  rW  ) r$   rŠ  rª  Úr_varsÚpw_varsÚ	red_numelr{  Ú
load_indexrM   rE  rP  Úactual_stridesr8  Ústrides_or_noneÚr_onlyÚpvÚtermrW  r&   Úmatchedr¦  ÚinnerÚstartÚaxesÚr_coeffÚpw_coeffÚ	pw_striderñ  rî  rö  r÷  r|  s                               @@@@@r)   Ú_get_reduction_axesz PallasKernel._get_reduction_axesË  s“  ü€ ð ×$Ò$Øà $× 5Ñ 5× ;Ñ ;Ó =×P™˜˜1ÀÇÃ’!ÐPˆÑPØ!%×!6Ñ!6×!<Ñ!<Ó!>×U™˜˜AÀaÇnÃn’1ÐUˆÑUÙ™WØà×1Ñ1Ó3ˆ	Ù˜I¨šNØà$(×$9Ñ$9×$?Ñ$?Ó$Aó A	/Ñ ˆH�jØ×(Ñ(¨Ó2ˆDØˆ|ØØ04Ñ-ˆAˆx˜˜N¨AÜ�X“ˆBØ�AŠvØØ:HÖI°Q˜tŸ~™~¨aÕ0ÐIˆOÐIÜÑ6 oÔ6Ô6ØÜ!%¤d¬3¡i°Ó!AˆGð  ˆFØò ,�ØŸ™ R¨Ó+‘ð,ä(2«ˆHÜŸ	™	×+Ñ+¨FÓ3ò $�Ø—>’>ØØ×,Ñ,Ó.‘��qØ—N‘N 5Ó)�Ø‘= Q¨£UØ—L‘L •Oð$ñ Ùô Ü@œE "›IÔ@óˆGñ Ùô
 �8‹}˜qÒ Ü�x“=¤C¨£LÒ0´S¸³[À1²_Ü Ó 9°Ô 9Ó9Ò9Ùô œD ›NÓ+ˆHØ ‘M XÑ-ˆDÜÓSÄUÈ2ÐPQÉ6Ã]ÔSÓSˆMÙà ™�Ø�Ø˜a’i D¨7°5¸1±9Ñ+=Ò$=Ø˜Q‘J�Eð ˜a’i D¨7°5¸1±9Ñ+=Ó$=äœE %¨°©Ó3Ó4‘ô õ ä" 2›Yôó �ñ
 Ü ›=�DÜÓ.¨Ô.Ó.Ò.ðCA	/ôH œ$˜t×4Ñ4×;Ñ;Ó=Ó>Ó?ˆ
Ø×"Ñ" 6¨!¡9Ó-ˆØ.5¸ªl�4—>‘> 'Ô*ÀˆØÐØˆHØ×#Ñ# G¨A¡JÓ/ˆØ08¸A²�D—N‘N 8Ô,À1ˆ	ØÐØˆIØ�iÒØØùóq QùÛUùò  Js   °NÁNÁ)NÁ>NÃ4N#c                ó„   — d}|D ]8  }|| j                   v sŒ| j                  | j                   |   «      }|€ y||z  }Œ: |S )zBCompute total numel for given prefixes (e.g., pointwise prefixes).r   N)ÚnumelsrG  )r$   Úprefixesr1   r'   Únumels        r)   Ú_compute_prefix_numelz"PallasKernel._compute_prefix_numel1  sP   € àˆØò 	 ˆAØ�D—K‘KÒØŸ™ t§{¡{°1¡~Ó6�Ø�=ÙØ˜%‘‘ð	 ð ˆr+   c                óŽ   — d}| j                   D ]3  }|j                  sŒ| j                  |j                  «      }|€ y||z  }Œ5 |S )zCompute total reduction numel.r   N)r  r   rG  r  )r$   r1   Útreer  s       r)   rû  z%PallasKernel._compute_reduction_numel<  sO   € àˆØ×$Ñ$ò 	 ˆDØ× Ó ØŸ™ t§z¡zÓ2�Ø�=ÙØ˜%‘‘ð	 ð ˆr+   c                ód  ‡ — ‰ j                   ry‰ j                  «       }|�|dkD  ryg }‰ j                  j                  D ]x  }‰ j	                  |«      }|€ y|\  }}}}}	|	s y t        |dd„ «      «       }
|
�ddl}|
|j                  k(  r yt        ˆ fd„|D «       «      }d|v r y|j                  |«       Œz |rt        t        |«      «      dkD  ry|r[d}|d   D ]  }||z  }Œ	 d}‰ j                  D ]3  }|j                  rŒ‰ j                  |j                  «      }|€ y||z  }Œ5 ||k7  ryy)	aŽ  
        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)
        FNr   Ú	get_dtypec                  ó   — y rS   r?   r?   r+   r)   r„  z4PallasKernel._can_use_tma_approach.<locals>.<lambda>g  r  r+   r   c              3  ó@   •K  — | ]  }‰j                  |«      –— Œ y ­wrS   ©rG  ©r@  r8  r$   s     €r)   rA  z5PallasKernel._can_use_tma_approach.<locals>.<genexpr>o  s   øè ø€ ÒD°a §¡¨q× 1ÑDùs   ƒT)rô  rû  r#   Úinput_buffersrF  r…  rµ   Úfloat64r¾  rý  r  r   r  r   rG  r  )r$   Úreduction_numelÚinput_shapesr£  rM   Úbuf_objrP  Ú	buf_numelr  r¦  Ú	buf_dtyperµ   Úshape_tupleÚinput_numelr8  ra  r  r  s   `                 r)   Ú_can_use_tma_approachz"PallasKernel._can_use_tma_approachG  sq  ø€ ð ×&Ò&Øð ×7Ñ7Ó9ˆØÐ&¨?¸QÒ+>Øð %'ˆØ—I‘I×+Ñ+ò 	-ˆDØ×(Ñ(¨Ó.ˆDØˆ|ÙØJNÑGˆG�X˜y¨.¸-Ù Ùð
 Dœ ¨±lÓCÓEˆIØÐ$Ûà §¡Ò-Ù ô  ÓD¸8ÔDÓDˆKØ�{Ñ"ÙØ×Ñ Õ,ð-	-ñ2 œC¤
¨<Ó 8Ó9¸AÒ=Øñ ØˆKØ! !‘_ò !�Ø˜qÑ ‘ð!ð ˆLØ×(Ñ(ò *�Ø×(Ó(Ø ŸN™N¨4¯:©:Ó6�EØ�}Ù$Ø  EÑ)‘Lð*ð ˜{Ò*Øàr+   c                óv  — t         j                  j                  |«      }|€y|j                  «       }d}|D ]  }| j	                  |«      }||�|n|z  }Œ g }d} t        |dd„ «      «       }	|	rt        |	dd«      nd}
|
�´t        t        |«      «      D ]'  }| j	                  |
|   «      }|j                  |«       Œ) t        |«      dk(  r|d   �c|d   dk7  r[d}nXt        |«      dkD  rJd}t        t        |«      dz
  d	d	«      D ],  }||   }|�||k7  rd}| j	                  ||   «      }|€Œ(||z  }Œ. |||||fS )
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                  ó   — y rS   r?   r?   r+   r)   r„  z/PallasKernel._get_buffer_info.<locals>.<lambda>ž  r  r+   r&  r   Fr·  )	r   r·   r¢  r£  rG  r…  rK  r  rý  )r$   r£  r%  rP  r&  r8  Úsvalr  r¦  rˆ  Ú
buf_striderÀ  Úactual_strideÚexpected_strideÚdim_sizes                  r)   rF  zPallasKernel._get_buffer_infoŒ  s}  € ô
 —'‘'×$Ñ$ TÓ*ˆØˆ?ØØ×#Ñ#Ó%ˆØˆ	Øò 	9ˆAØ—>‘> !Ó$ˆDØ Ð!1™°qÑ8‰Ið	9ð
  "ˆØˆà=”˜ ,±Ó=Ó?ˆÙ8>”W˜V X¨tÔ4ÀDˆ
àÐ!Üœ3˜x›=Ó)ò 5�Ø $§¡¨z¸!©}Ó =�Ø×%Ñ% mÕ4ð5ô
 �8‹} Ò!Ø! !Ñ$Ð0°^ÀAÑ5FÈ!Ò5KØ$)‘MÜ�X“ Ò"Ø"#�Üœs 8›}¨qÑ0°"°bÓ9ò 4�AØ$2°1Ñ$5�MØ$Ð,°ÀÒ0PØ(-˜Ø#Ÿ~™~¨h°q©kÓ:�HØÑ+Ø'¨8Ñ3™ð4ð ˜ )¨^¸]ÐJÐJr+   c                óô   — | j                  |«      }g }|D ]O  }|| j                  v sŒ| j                  |   }| j                  |j                  «      }|€Œ?|j	                  |«       ŒQ d}|D ]  }||z  }Œ	 ||fS )zNCompute expected output numel and used vars from iteration variables in index.r   )rÊ   r/  rG  rO  rý  )	r$   r  r$  Úused_range_lengthsrÔ   rÏ  rÑ  ra  Úls	            r)   rM  z-PallasKernel._compute_output_numel_from_index¶  sš   € ð ×,Ñ,¨UÓ3ˆ	àÐØò 	:ˆCØ�d×+Ñ+Ò+Ø×-Ñ-¨cÑ2�Ø!Ÿ^™^¨E¯L©LÓ9�
ØÑ)Ø&×-Ñ-¨jÕ9ð	:ð ˆØ#ò 	ˆAØ˜AÑ‰Lð	ð ˜YÐ&Ð&r+   c                óÔ   — t        «       }|D ]X  }t        j                  ||«      }t        j                  ||«      }|€d}| j	                  |«      }|j                  |�|n|«       ŒZ |S )zT
        Extract coefficients of iteration variables from index expression.
        r   )r   r   r   r!  rG  r  )r$   r  r$  ÚcoefficientsrÔ   r%  r&  Úcoefs           r)   Ú_get_index_coefficientsz$PallasKernel._get_index_coefficientsÊ  su   € ô $.£<ˆØò 	CˆCÜ*×GÑGÈÈsÓSˆHÜ(×@Ñ@ÀÈ3ÓOˆFØˆ~Ø�Ø—>‘> &Ó)ˆDØ×Ñ TÐ%5™T¸6ÕBð	Cð Ðr+   c                ó:  — dg}t        |«      dkD  rPd}g }t        t        |«      dz
  dd«      D ]0  }|j                  d|«       | j                  ||   «      }|€Œ,||z  }Œ2 |rt	        |«      }	|D ]  }
|
|	vsŒ y yt	        d„ |D «       «      }|D ]  }
|
|vsŒ y y)zR
        Check if access pattern requires gather (non-standard striding).
        r   r·  r   Tc              3  ó&   K  — | ]	  }|€Œ|–— Œ y ­wrS   r?   r?  s     r)   rA  z5PallasKernel._check_gather_pattern.<locals>.<genexpr>ø  s   è ø€ Ò*V°ÈÉ¬1Ñ*Vùs   ‚ŠF)r  rK  ÚinsertrG  r   )r$   rP  r  r¦  r6  Úexpected_stridesr0  rÀ  r1  Úexpected_stride_setr7  Úactual_stride_sets               r)   Ú_check_gather_patternz"PallasKernel._check_gather_patternÚ  sÚ   € ð ˜3Ðäˆx‹=˜1ÒØˆOØ!ÐÜœ3˜x›=¨1Ñ,¨b°"Ó5ò 0�Ø ×'Ñ'¨¨?Ô;ØŸ>™>¨(°1©+Ó6�ØÑ'Ø# xÑ/‘Oð	0ñ ä",Ð-=Ó">ÐØ$ò  �ØÐ2Ò2Ùð ð ô !+Ñ*V°nÔ*VÓ VÐØ$ò  �ØÐ0Ò0Ùð ð r+   c                ó  ‡ — |j                   dk7  s|j                  r|S ‰ j                  |«      }|€|S |\  }}}}}	‰ j                  |«      \  }
}‰ j	                  «       }‰ j                  ||«      }‰ j                  |||	|«      }t        ˆ fd„|D «       «      }t        |«      t        |«      k  xr. t        |«      dkD  xr |dkD  xr t        |«      t        |«      kD  }t        j                  j                  «       j                  dk(  }|	 xr t        d„ |D «       «      }t        d„ |D «       «      }|xr
 | xr ||
k(  }|
dkD  r7||
k7  s|s|r.t        |«      dkD  r |s|st        ‰ j!                  |«      d¬	«      S |S )
zVCheck if buffer access needs strided indexing due to size mismatch or gather patterns.r  c              3  óL   •K  — | ]  }‰j                  |«      d k7  sŒd –— Œ y­wrD  r  r   s     €r)   rA  z7PallasKernel._needs_strided_indexing.<locals>.<genexpr>  s!   øè ø€ Ò O q¸¿¹ÀqÓ8IÈQÓ8N¤Ñ Oùs   ƒ$�$r   r   r³   c              3  ó$   K  — | ]  }|d u–— Œ
 y ­wrS   r?   r?  s     r)   rA  z7PallasKernel._needs_strided_indexing.<locals>.<genexpr>#  s   è ø€ ò <
ØˆA�TŒMñ<
ùrš  c              3  óL   K  — | ]  }t        |t        t        z  «       –— Œ y ­wrS   )r§   rÀ  r¨   r�  s     r)   rA  z7PallasKernel._needs_strided_indexing.<locals>.<genexpr>&  s   è ø€ ÒUÀ1¤J¨q´#¼±+Ó$>Ô >ÑUùs   ‚"$TrÕ  )rÈ  rÉ  rF  rM  r+  r8  r?  Úsumr  r   r·   r¸   r¹   rL  rH  rÇ  r-  )r$   r£  r  ÚindexingrM   r%  rP  r&  r  r¦  ra  r$  Úall_iter_varsr6  Úhas_non_unit_stridesÚbuf_effective_dimsÚnot_all_vars_usedrÖ  Úis_known_non_contiguousÚhas_symbolic_coefÚskip_for_non_contiguouss   `                    r)   Ú_needs_strided_indexingz$PallasKernel._needs_strided_indexingÿ  s´  ø€ ð ×Ñ Ò&¨(×*@Ò*@ØˆOà×$Ñ$ TÓ*ˆØˆ<ØˆOàFJÑCˆ�˜9 n°mØ"&×"GÑ"GÈÓ"NÑˆ�iØ×+Ñ+Ó-ˆØ×3Ñ3°E¸9ÓEˆð  $×9Ñ9Ø�n 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ÀLÑ7Pð 	 ð ˜1ÒØ˜lÒ*Ñ.?ÑCWÜ�I“ Ò"Ù+Ù%ä"Ø×6Ñ6°uÓ=ÈTôð ð ˆr+   c                óœ  — |j                   s|j                  dk(  r|S t        j                  j	                  |«      }|€|S |j                  «       }t        |«      dk(  rt        d|j                   ¬«      S t        |«      dk(  r0|d   dk(  r(| j                  |«      st        d|j                   ¬«      S t        |«      dkD  rT| j                  |«      }|st        |j                  d¬«      S d|j                  v rt        | j                  |«      d¬«      S | j                  r*d|j                  v rt        | j                  |«      d¬«      S |S )z`
        Adjust index expression based on buffer shape (0-dim scalar, multi-dim, etc.).
        r  r   rÕ  r   TrÖ  )rÉ  rÈ  r   r·   r¢  r£  r  rÇ  r4  r-  rï  )r$   r£  r  rE  r%  rP  rÜ  s          r)   Ú_adjust_index_for_buffer_shapez+PallasKernel._adjust_index_for_buffer_shape9  sP  € ð ×!Ò! X×%7Ñ%7¸5Ò%@ØˆOä—'‘'×$Ñ$ TÓ*ˆØˆ?ØˆOà×#Ñ#Ó%ˆäˆx‹=˜AÒÜ"Ø¨x×/EÑ/Eôð ô �‹M˜QÒØ˜‘˜qÒ Ø×,Ñ,¨UÔ3ä"Ø¨x×/EÑ/Eôð ô
 ˆx‹=˜1ÒØ ×4Ñ4°UÓ;ˆMÙ Ü&Ø&×0Ñ0Àôð ð ˜×+Ñ+Ñ+ä&Ø"×:Ñ:¸5ÓAÐQUôð ð
 �;Š;˜4 8×#5Ñ#5Ñ5Ü"Ø×6Ñ6°uÓ=ÈTôð ð ˆr+   c                ó¼  — |j                   s|S t        j                  j                  |«      }|€|S |j	                  «       }t        |«      }|dk  r|S | j                  |«      }t        |«      dk7  r|S t        t        |«      «      }t        j                  ||«      }	| j                  t        j                  |	|«      «      }
|
�|
dk  r|S t        j                  j                  j                  ||	z
  «      }	 t        |«      }|dk  s||
k\  r|S | j                  |d   «      }|�||
z  dk7  r|S | j$                  j'                  |«      }|€|S | j                  |j(                  «      }d}|D ]  }| j                  |«      }|€|c S ||z  }Œ  |�||
z  |k7  r|S d|dz
  z  }|dk(  r|› d|
› �}n	|› |› d|
› �}t+        |d¬«      S # t         t"        f$ r |cY S w xY w)	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Ö  FrÕ  )rÉ  r   r·   r¢  r£  r  rÊ   r  r  r   r   rG  r!  r  r  rÀ  r"  r#  r/  r»  rO  rÇ  )r$   r£  r  rE  r%  rP  rS  r$  rÔ   r%  r&  r¤  r'  Úlast_dimrÏ  Ú
var_lengthr&  r8  rZ  r�  Ú	slice_strs                        r)   Ú_try_multidim_slicez PallasKernel._try_multidim_sliceo  s  € ð ×%Ò%ØˆOä—'‘'×$Ñ$ TÓ*ˆØˆ?ØˆOà×#Ñ#Ó%ˆÜ�8‹}ˆØ�!Š8ØˆOð ×,Ñ,¨UÓ3ˆ	Üˆy‹>˜QÒØˆOä”4˜	“?Ó#ˆÜ&×CÑCÀEÈ3ÓOˆØ—‘Ü×7Ñ7¸À#ÓFó
ˆð ˆ>˜V qš[ØˆOä—‘×!Ñ!×*Ñ*¨5°8Ñ+;Ó<ˆð	Ü˜V›ˆJð ˜Š>˜Z¨6Ò1ØˆOà—>‘> (¨2¡,Ó/ˆØÐ˜x¨&Ñ0°AÒ5ØˆOð ×%Ñ%×)Ñ)¨#Ó.ˆØˆ=ØˆOØ—^‘^ E§L¡LÓ1ˆ
Øˆ	Øò 	ˆAØ—‘˜qÓ!ˆAØˆyØ’Ø˜‰N‰Ið		ð
 Ð ¨fÑ!4¸	Ò!AØˆOà˜$ ™(Ñ#ˆØ˜Š?Ø!˜( " V HÐ-‰Ià!˜( : ,¨b°°Ð9ˆIÜ¨À%ÔHÐHøôA œ:Ð&ò 	ØŠOð	ús   Ã:G ÇGÇGc                ó   — d| › d|› d�S )a  Generate gather-based permutation instead of jnp.permute_dims.

        Avoids a Mosaic compiler bug where jnp.permute_dims produces
        corrupted output tensors on TPU for 3D+ arrays.  Uses
        pallas_permute which flattens to 1D and does a 1D gather.
        zpallas_permute(r    r!   r?   )Ú	load_exprr¯  s     r)   Ú_gather_permute_exprz!PallasKernel._gather_permute_expr¹  s   € ð !  ¨2¨d¨V°1Ð5Ð5r+   c                ó^  — || j                   v r| j                   |   S | j                  j                  D ]w  }t        |«      j	                  «       }|j                  |› d�«      sŒ1t        j                  d|j                  dd«      d   «      D ]  }| j                  |«      }|€Œ|c c S  Œy y)záTrace a tmp variable back to its source buffer's kernel param.

        Follows CSE assignments backward through bounds-checking (where/clamp)
        until it finds a variable that was directly loaded from a buffer.
        rr  ú
\btmp\d+\br   N)
r  rÐ   Ú_linesr:   rÙ  r7  ÚreÚfindallÚsplitÚ_trace_to_load_source)r$   Úvar_nameÚlineÚline_strÚrefr1   s         r)   r^  z"PallasKernel._trace_to_load_sourceÃ  s°   € ð �t×)Ñ)Ñ)Ø×%Ñ% hÑ/Ð/Ø—L‘L×'Ñ'ò 	"ˆDÜ˜4“y×'Ñ'Ó)ˆHØ×&Ñ&¨(¨°3Ð'7Ô8ØÜ—z‘z -°·±ÀÀqÓ1IÈ!Ñ1LÓMò "�Ø×3Ñ3°CÓ8�ØÑ%Ø!”Mñ"ð		"ð r+   c                óô  ‡ ‡‡— ‰ j                  |«      }|€y|\  }}}}}|D �cg c]  }‰ j                  |«      ‘Œ }}t        |«      dk  st        d„ |D «       «      ryt	        t
        t           |«      }	‰ j                  |«      }
t        |
«      dk7  ry|
d   }‰ j                  ||«      }|dk(  st        |t        «      sy‰ j                  |	«      }‰ j                  |g|«      }|€y|d   }t        |	«      }|t        d|dz
  «      k\  ryt        |«      }‰ j                  |«      }|€y‰ j                  j!                  |«      }|�’‰ j                  |«      }|�|\  }}}}}t        |«      dk7  ry‰ j                  |d   «      €yt#        j$                  ˆ ˆfd„|D «       «      }t#        j$                  ˆˆ fd„‰ j'                  |«      D «       «      }||k\  ryt)        ||t+        |	«      |||¬«      S c c}w )	a,  Detect a load with data-dependent indexing suitable for scalar prefetch.

        Matches exactly one indirect variable whose coefficient corresponds to
        a C-contiguous stride dimension.  Rejects 1-to-1 gather patterns where
        the indices buffer covers the full iteration space.
        Nr   c              3  ó$   K  — | ]  }|d u –— Œ
 y ­wrS   r?   r?  s     r)   rA  z7PallasKernel._detect_indirect_access.<locals>.<genexpr>ã  s   è ø€ Ò'H°a¨¨T¬	Ñ'Hùrš  r   r   c              3  óJ   •K  — | ]  }‰j                  |«      xŠ€Œ‰–— Œ y ­wrS   r  )r@  r8  r$   rŠ  s     €€r)   rA  z7PallasKernel._detect_indirect_access.<locals>.<genexpr>
	  s)   øè ø€ ò *Ø°T·^±^ÀAÓ5FÐ0F°Ñ/S”Añ*ùs   ƒ#œ#c              3  ó”   •K  — | ]?  }|‰j                   v r/‰j                  ‰j                   |   j                  «      xŠ	 �‰–— ŒA y ­wrS   )r/  rG  rO  )r@  rÔ   rO  r$   s     €€r)   rA  z7PallasKernel._detect_indirect_access.<locals>.<genexpr>	  sQ   øè ø€ ò )àØ˜d×3Ñ3Ñ3Ø"&§.¡.°×1FÑ1FÀsÑ1K×1RÑ1RÓ"SÐS˜Øð ô ñ)ùs   ƒAA)r¼  r½  r¿  rÁ  rÂ  rÃ  )rF  rG  r  rH  r   rI  rÀ  r9  r¥  r§   rJ  r¤  Úmaxr:   r^  r  r»  r©   r¡  rÊ   r»  r¾  )r$   rl  r£  r  Úbuf_inforE  rP  r8  Úbuf_size_rawÚbuf_size_intsÚindirect_varsrÂ  rW  r|  rã  rÁ  rS  Úindirect_var_namerÃ  Úindices_graph_nameÚindices_infoÚindices_sizeÚindices_numelÚiter_productrO  rŠ  s   `                       @@r)   Ú_detect_indirect_accessz$PallasKernel._detect_indirect_accessÕ  s(  ú€ ð ×(Ñ(¨Ó.ˆØÐØØ'Ñˆˆ8�Q˜˜1Ø3;Ö<¨a˜Ÿ™ qÕ)Ð<ˆÐ<Üˆ|Ó˜qÒ ¤CÑ'H¸<Ô'HÔ$HØÜ#'¬¬S©	°<Ó#@ˆà×/Ñ/°Ó6ˆÜˆ}Ó Ò"ØØ$ QÑ'ˆà×+Ñ+¨E°<Ó@ˆØ�AŠ:œZ¨¬sÔ3Øð ×,Ñ,¨]Ó;ˆØ×*Ñ*¨E¨7°GÓ<ˆØˆ?ØØ˜q‘zˆä�=Ó!ˆØœ3˜q $¨¡(Ó+Ò+Øä Ó-ÐØ×2Ñ2Ð3DÓEˆØÐ Øð "×6Ñ6×:Ñ:¸=ÓIÐØÐ)Ø×0Ñ0Ð1CÓDˆLØÐ'Ø+7Ñ(��<  A qÜ�|Ó$¨Ò)ØØ—>‘> ,¨q¡/Ó2Ð:ØÜ $§	¡	ô *Ø+ô*ó !�ô  $Ÿy™yô )à#×7Ñ7¸Ó>ô)ó  �ð ! LÒ0Øä"ØØÜ˜mÓ,Ø%Ø*Ø'ô
ð 	
ùòk =s   ¤G5c                óL  — t        «       }| j                  D ]?  \  }}t        j                  d|«      D ]!  }|j	                  |j                  «       «       Œ# ŒA g }| j                  j                  D ]‰  }t        |«      j                  «       }t        j                  d|t        j                  «      }|r3|j                  |j                  d«      |j                  d«      |f«       Œv|j                  d||f«       Œ‹ d}|rpd}t        |«      D ]]  \  }	}
}|	sŒ
|	|v sŒt        j                  d|
«      D ]6  }|j                  «       |vsŒ|j	                  |j                  «       «       d}Œ8 Œ_ |rŒp|D �	��cg c]  \  }	}}|	�|	|v r|‘Œ c}}}	| j                  _        yc c}}}	w )aj  Remove dead compute lines after scalar prefetch replaces indirect load.

        When the table load is simplified to buf[0] (scalar prefetch handles
        indexing), the indices load and all derived bounds-checking code become
        dead.  This performs backward liveness analysis from the store variables
        to identify and remove dead lines.
        rY  z^(tmp\d+)\s*=\s*(.*)r   r   NTF)r   rñ  r[  Úfinditerr  ÚgrouprÐ   rZ  r:   rÙ  ÚmatchÚDOTALLrý  rŽ  )r$   Ú	live_varsrE  Ú
store_lineÚmÚassignmentsr`  ra  Úchangedr_  Úrhss              r)   Ú_eliminate_dead_indirect_codez*PallasKernel._eliminate_dead_indirect_code 	  sŒ  € ô &0£\ˆ	Ø!×3Ñ3ò 	)‰MˆAˆzÜ—[‘[ °
Ó;ò )�Ø—‘˜aŸg™g›iÕ(ñ)ð	)ð
 :<ˆØ—L‘L×'Ñ'ò 	;ˆDÜ˜4“y×'Ñ'Ó)ˆHÜ—‘Ð0°(¼B¿I¹IÓFˆAÙØ×"Ñ" A§G¡G¨A£J°·±¸³
¸DÐ#AÕBà×"Ñ" D¨(°DÐ#9Õ:ð	;ð ˆÙØˆGÜ$,¨[Ó$9ò +Ñ �˜#˜qÚ ¨IÒ 5ÜŸ[™[¨¸Ó<ò +˜ØŸ7™7›9¨IÒ5Ø%ŸM™M¨!¯'©'«)Ô4Ø&*™Gñ+ð+ò ð &1÷
ð 
á!�˜!˜TØÐ 8¨yÑ#8ò ô
ˆ�‰Õùô 
s   Å7Fc                ó(  ‡— |j                   rÖ| j                  |||«      }|�.| j                  �|| j                  k(  sJ d«       ‚|| _        |› d�S d| _        | j                  j                  |«       |j                  t        j                  t        j                  t        t        «      }| j                  rdnd}|rd|j                  › d|› d�n|j                  }|› d	|› d
�S |› d|j                  › d
�}	|j                  dk(  r�| j                  s�| j                  ||«      }
|
�#| j!                  |	|
«      }	|
| j"                  |<   |	S | j%                  ||«      }|�6|\  Š}d|	› d|› d�}	‰| j&                  |<   t)        ˆfd„|D «       «      | _        |	S )zC
        Build the load expression based on indexing mode.
        z-only one indirect access per kernel supportedz[0]Tz	jnp.int32z	jnp.int64ré   r´   r!   z[...].flatten()[rk  rj  r  zjnp.permute_dims(r    c              3  ó(   •K  — | ]	  }‰|   –— Œ y ­wrS   r?   )r@  r'   Úcollapsed_shapes     €r)   rA  z0PallasKernel._build_load_expr.<locals>.<genexpr>~	  s   øè ø€ ò <Ø34˜O¨AÕ.ñ<ùrò  )rÉ  rr  r  rÿ  r  r  r  r  r	   r   ÚTorchMinÚTorchMaxrÖ  rÈ  rï  rµ  rW  rô  rÌ  rõ  r¾  rö  )r$   rl  r£  r  rE  ÚindirectÚ
has_minmaxÚ	idx_dtyperÌ  rV  r¯  Ú	collapsedÚcpermr�  s                @r)   Ú_build_load_exprzPallasKernel._build_load_exprJ	  sÑ  ø€ ð ×!Ò!à×3Ñ3°C¸¸uÓEˆHØÐ#Ø×'Ñ'Ð3ð $ t×';Ñ';Ò;ð ØGóÐ;ð (0�Ô$Ø˜˜c�{Ð"à(,ˆDÔ%Ø×(Ñ(×,Ñ,¨TÔ2àŸ™¤5§9¡9¬e¯i©i¼Ä8ÓLˆJØ'+§{¢{™¸ˆIñ ð �H×&Ñ&Ð' y°°¸1Ñ=à×'Ñ'ð ð
 �UÐ*¨3¨%¨qÐ1Ð1ð ˜%˜q ×!3Ñ!3Ð 4°AÐ6ˆIà×!Ñ! UÒ*°4·;²;Ø×6Ñ6°t¸UÓC�ØÐ#Ø $× 9Ñ 9¸)ÀTÓ J�IØ8<�D×/Ñ/°Ñ5ð Ðð !%× DÑ DÀTÈ5Ó Q�IØ Ð,Ø1:Ñ.˜¨Ø&7¸	°{À"ÀUÀGÈ1Ð$M˜	ð ?N˜×5Ñ5°dÑ;Ü6;ó <Ø8=ô<ó 7˜Ô3ð Ðr+   c                ó  ‡— |j                  d«      s|S t        ˆfd„| j                  j                  D «       «      }|rMt        j
                  j                  |«      Š‰�,‰j                  «       }t        |«      dk(  r|d   dk(  rd|› d�S |S )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,)
        rl  c              3  óÈ   •K  — | ]Y  }|j                  d «       xrA t        j                  j                  |«      xŠduxr t	        ‰j                  «       «      dk(  –— Œ[ y­w)rl  Nr   )r7  r   r·   r¢  r  r£  )r@  r{  r%  s     €r)   rA  zBPallasKernel._maybe_squeeze_intermediate_buffer.<locals>.<genexpr>Ž	  sj   øè ø€ ò 
ð ð ×#Ñ# EÓ*Ð*ò -ÜŸG™G×.Ñ.¨xÓ8Ð8�ÀÐEò-ä�G×$Ñ$Ó&Ó'¨1Ñ,ó-ñ
ùs   ƒAA"r   r·  r   zjnp.squeeze(z
, axis=-1))	r7  rH  r#   r!  r   r·   r¢  r£  r  )r$   r£  rV  Úhas_1d_inputrP  r%  s        @r)   Ú"_maybe_squeeze_intermediate_bufferz/PallasKernel._maybe_squeeze_intermediate_buffer„	  s•   ø€ ð �‰˜uÔ%ØÐô ó 
ð !ŸI™I×3Ñ3ô	
ó 
ˆñ Ü—g‘g×(Ñ(¨Ó.ˆGØÐ"Ø"×+Ñ+Ó-�Ü�x“= AÒ%¨(°2©,¸!Ò*;Ø)¨)¨°JÐ?Ð?àÐr+   c                ó  — t         j                  j                  |«      }|�t        |j	                  «       «      dk7  r|S |j                  d«      r|S | j                  |j	                  «       d   «      }|€|S t         j                  j                  |«      }|�|j                  s|S d}| j                  j                  D ]ƒ  }t         j                  j                  |«      }	|	€Œ%t        |	j	                  «       «      dkD  sŒB|	j	                  «       D �
cg c]  }
| j                  |
«      ‘Œ }}
t        d„ |D «       «      r nd}Œ… |�t        |«      dk  r|S | j                  |«      }t        |«      dk7  r|S t        t        |«      «      }|| j                  vr|S | j                  |   }| j                  |j                   «      |k7  r|S | j                  j#                  «       D ��cg c]1  \  }}| j                  |j                   «      |k(  r|j$                  s|‘Œ3 }}}t        |«      dk7  r|S d}t'        |«      D ��
cg c]  \  }}
|
|k(  sŒ|‘Œ }}}
t        |«      dk(  r|d   }n?d}| j                  j#                  «       D ]   \  }}||k(  r|} n|j$                  rŒ|dz  }Œ" |€|S |t        |«      dz
  k(  r|S dgt        |«      z  }d||<   |› ddj)                  t+        t,        |«      «      › d	�S c c}
w c c}}w c c}
}w )
aÂ  Reshape 1D buffers for higher-dim broadcasting in reduction kernels.

        When a 1D buffer (e.g. a reduction result from a prior kernel, or a
        batch-norm parameter) is loaded into a kernel with 2+ iteration dims,
        JAX right-aligns it for broadcasting: (N,) becomes (1, N).  This is
        wrong when the buffer corresponds to a non-trailing axis; we reshape
        to (N, 1, ...) so broadcasting matches the correct axis.
        Nr   rl  r   c              3  ó$   K  — | ]  }|d u–— Œ
 y ­wrS   r?   r?  s     r)   rA  z:PallasKernel._maybe_broadcast_1d_buffer.<locals>.<genexpr>Á	  s   è ø€ Ò;¨�q ”}Ñ;ùrš  r·  rs  r    r!   )r   r·   r¢  r  r£  r7  rG  r  Úis_floating_pointr#   r!  rL  rÊ   r  r  r/  rO  rN  r   rw  rV  Úmapr:   )r$   r£  r  rV  r%  Ú
buf_lengthr»   Úref_buf_sizer{  Ú	other_bufr8  r$  Úused_varrÏ  rŠ  rª  Úmatching_varsÚaxis_posrÀ  Úmatching_dimsÚpw_idxÚsymÚreshape_dimss                          r)   Ú_maybe_broadcast_1d_bufferz'PallasKernel._maybe_broadcast_1d_bufferž	  s  € ô —'‘'×$Ñ$ TÓ*ˆØˆ?œc '×"2Ñ"2Ó"4Ó5¸Ò:ØÐð
 �?‰?˜5Ô!ØÐà—^‘^ G×$4Ñ$4Ó$6°qÑ$9Ó:ˆ
ØÐØÐä—‘×!Ñ! $Ó'ˆØÐ U×%<Ò%<ØÐð ˆØŸ	™	×/Ñ/ò 	$ˆHÜŸ™×*Ñ*¨8Ó4ˆIØÑ$¬¨Y×-?Ñ-?Ó-AÓ)BÀQÓ)FØ;D×;MÑ;MÓ;OÖP°a §¡¨qÕ 1ÐP�ÐPÜÑ;¨lÔ;Ô;ÙØ#‘ð	$ð Ð¤3 |Ó#4¸Ò#9ØÐð ×,Ñ,¨UÓ3ˆ	Üˆy‹>˜QÒØÐÜœ˜Y›Ó(ˆØ˜4×0Ñ0Ñ0ØÐð ×%Ñ% hÑ/ˆØ�>‰>˜%Ÿ,™,Ó'¨:Ò5ØÐð ×-Ñ-×3Ñ3Ó5÷
á��1Ø�~‰~˜aŸh™hÓ'¨:Ò5¸a¿nºnò ð
ˆñ 
ô
 ˆ}Ó Ò"ØÐð ˆÜ'0°Ó'>×R™t˜q !À!ÀzÃ/šÐRˆÑRÜˆ}Ó Ò"Ø$ QÑ'‰Hð ˆFØ×/Ñ/×5Ñ5Ó7ò  ‘��QØ˜(’?Ø%�HÙØ—~“~Ø˜a‘K‘Fð ð ÐØÐØ”s˜<Ó(¨1Ñ,Ò,ØÐà�sœS Ó.Ñ.ˆØ!#ˆ�XÑØ�˜I d§i¡i´´C¸Ó0FÓ&GÐ%HÈÐJÐJùòo  Qùó.
ùó Ss   ÄK:Ç)6K?ÉLÉLc                ó„  — |j                   dk7  s|j                  r|S | j                  |«      }| j                  «       }| j	                  |«      }t        |d«      r|j                  n	t        «       |z  }||z
  }|rt        |«      dk  r|S d}| j                  j                  «       D ]w  \  }	}
| j	                  |
«      }|sŒ||k7  rŒ | j                  |
«      }t        |d«      r|j                  n	t        «       |z  }||k7  s||k(  rŒa| j                  |	|
|«      rŒud} n |rt        | j                  |«      d¬«      S |S )a
  
        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  r*  r   FTrÕ  )rÈ  rÉ  rË   r+  rÊ   Úhasattrr*  r   r  rò  rN  Ú_check_load_is_strided_inputrÇ  r-  )r$   r  rE  r  r,  Ústore_orig_varsÚstore_prep_varsÚnew_varsÚhas_im2col_patternr{  r  Úload_orig_varsÚ	prep_loadÚload_prep_varss                 r)   Ú_check_im2col_patternz"PallasKernel._check_im2col_patternù	  s{  € ð ×Ñ Ò&¨(×*@Ò*@ØˆOà×.Ñ.¨uÓ5ˆØ×'Ñ'Ó)ˆ	Ø×2Ñ2°5Ó9ˆô �~ ~Ô6ð ×'Ò'ä“Øñ	ˆð
 # _Ñ4ˆñ œ3˜Ó/°1Ò4ØˆOð #ÐØ$(×$9Ñ$9×$?Ñ$?Ó$Aò 	Ñ ˆH�jØ!×5Ñ5°jÓAˆNÙ!Øð  Ò0Øð ×-Ñ-¨jÓ9ˆIô ˜9 nÔ5ð ×&Ò&ä“\Øñ	ˆNð  Ò/°?ÀoÒ3UØð ×4Ñ4Ø˜* nõð &*Ð"Ùð5	ñ8 Ü"Ø×6Ñ6°~ÓFØ"ôð ð
 ˆr+   c                óH  — t         j                  j                  |«      }|€y t        |dd„ «      «       }|€yt        |dd«      }|€y|j	                  «       }g }|D ]W  }	t        j                  ||	«      }
t        j                  |
|	«      }|€Œ2| j                  |«      }|j                  |�|n|«       ŒY t        «       }t        |«      D ]G  \  }}| j                  ||   «      }|�|dkD  sŒ"| j                  |«      }|j                  |�|n|«       ŒI t        |«      |k(  S )z\
        Check if load coefficients match buffer strides (strided input vs im2col).
        NFr‚  c                  ó   — y rS   r?   r?   r+   r)   r„  z;PallasKernel._check_load_is_strided_input.<locals>.<lambda>C
  r  r+   r&  r   )r   r·   r¢  r…  r£  r   r   r!  rG  rý  r   rw  r  )r$   r{  r  r¤  rl  rˆ  Úbuf_stridesÚ	buf_sizesÚload_coeffsrÔ   r%  r7  Úint_coefÚbuf_stride_setrÀ  r8  r1  Úint_ss                     r)   rŸ  z)PallasKernel._check_load_is_strided_input9
  s9  € ô �g‰g× Ñ  Ó*ˆØˆ;Øà9”˜˜l©LÓ9Ó;ˆØˆ>Øä˜f h°Ó5ˆØÐØà—L‘L“Nˆ	ð ˆØ!ò 	OˆCÜ*×GÑGÈ
ÐTWÓXˆHÜ&×>Ñ>¸xÈÓMˆDØÑØŸ>™>¨$Ó/�Ø×"Ñ"¨xÐ/C¡8ÈÕNð	Oô $›ˆÜ˜kÓ*ò 	F‰DˆAˆqØ—~‘~ i°¡lÓ3ˆHØÐ 8¨a£<ØŸ™ qÓ)�Ø×"Ñ"¨EÐ,=¡5À1ÕEð		Fô ˜+Ó&¨.Ñ8Ð8r+   c                óÞ  — | j                   ry| j                  |«      }|€y|\  }}}}}t        |«      dk7  st        |«      dk7  ry| j                  |d   «      }| j                  |d   «      }|d   }|d   }	|�|	�||	k  r|�|�
|dkD  r|dkD  sy| j                  j
                  D ]D  }
| j                  |
«      }|€Œ|\  }}}}}t        |«      dk7  rŒ.|d   }|d   }|€Œ;|€Œ>||k  sŒD y y)a  
        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)rô  rF  r  rG  r#   r!  )r$   r£  rM   rE  rP  r  Úsize0Úsize1Ús0Ús1Úinp_nameÚinp_infoÚinp_stridesÚinp_s0Úinp_s1s                  r)   Ú_check_store_needs_transposez)PallasKernel._check_store_needs_transposea
  s:  € ð ×&Ò&Øà×$Ñ$ TÓ*ˆØˆ<Øà,0Ñ)ˆˆ8�Q˜¨Üˆ~Ó !Ò#¤s¨8£}¸Ò'9Øà—‘˜x¨™{Ó+ˆØ—‘˜x¨™{Ó+ˆØ˜AÑˆØ˜AÑˆð ˆNØ�Ø�R’ØÐ!ØÐ!Ø˜’	Ø˜’	àð Ÿ	™	×/Ñ/ò 
	ˆHØ×,Ñ,¨XÓ6ˆHØÐØØ&.Ñ#ˆAˆq�!�[ !Ü�;Ó 1Ò$ØØ  ‘^ˆFØ  ‘^ˆFØÑ! fÑ&8¸VÀf»_Ùð
	ð r+   c                óŠ   — d|› d�g}|r|j                  |› d|› d�«       |S |j                  |› d|› d|› d|› d|› d�
«       |S )	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_exprz)PallasKernel._build_full_array_store_expr•
  s‹   € ð ' u g¨QÐ/Ð0ˆÙØ�L‰LØ�%ð Ø˜5ð !+ð,ôð ˆð �L‰LØ�%ð Ø˜5ð !&Ø&) UÐ*BÀ3À%ð H/Ø/2¨e°8ð=ôð ˆr+   c                ó,  — |j                   dk(  r$| j                  |«      }| j                  |||«      S |j                  r/d| _        |dk(  rdnd}|› d|› d|j                   › d|› d	|› d
|› d�gS | j                  |«      }	t        j                  j                  |«      }
|
�B|
j                  «       }t        |«      dkD  r$| j                  |«      s| j                  ||d«      S |	rœ|dk(  rdnd}d|› d�g}d|j                   › d|› d�}|dk(  rO| j                  j                  |«       |› d�}|j                  |› d|› d|j                   › d|› d	|› d
|› d�«       |S |j                  |› d|j                   › d|› �«       |S |› d|j                   › d|› �gS )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  TÚ
atomic_addr  Úsetú[...] = z[...].flatten().at[jnp.asarray(z).flatten()].z(jnp.asarray(z).flatten()).reshape(ú.shape)r   Fr¼  r!   z
(jnp.full(z%.shape, _val) if _val.ndim == 0 else Ú_aliasrj  z] = )rÈ  rº  rÂ  rÉ  rÿ  r;  r   r·   r¢  r£  r  r4  ró  r  rý  )r$   r¾  r£  r  r¿  rE  ÚmoderÀ  Ú
scatter_oprÛ  rl  rP  rÁ  r×   Úalias_params                  r)   Ú_build_store_exprzPallasKernel._build_store_expr®
  s  € ð ×Ñ Ò&à"×?Ñ?ÀÓEˆOØ×4Ñ4°S¸%ÀÓQÐQà×!Ò!Ø(,ˆDÔ%à"&¨,Ò"6™¸EˆJà�%�x ˜uÐ$CÀH×DVÑDVÐCWÐWdÐeoÐdpð qØ$˜gÐ%:¸3¸%¸wðHðð ð ×.Ñ.¨uÓ5ˆÜ�g‰g× Ñ  Ó&ˆàˆ?Ø—|‘|“~ˆHÜ�8‹}˜qÒ ¨×)AÑ)AÀ%Ô)Hà×8Ñ8¸¸eÀUÓKÐKáà"&¨,Ò"6™¸EˆJØ*¨5¨'°Ð3Ð4ˆEØ% h×&8Ñ&8Ð%9Ð9^Ð_dÐ^eÐefÐgˆJØ�|Ò#à×&Ñ&×*Ñ*¨3Ô/Ø!$  V˜n�Ø—‘Ø�e˜8 K =Ð0OÐPX×PbÑPbÐOcÐcpÐq{Ðp|ð }#Ø#- ,Ð.CÀCÀ5ÈðQôð ˆLð —‘ ˜u A h×&8Ñ&8Ð%9¸¸j¸\ÐJÔKØˆLà�%�q˜×+Ñ+Ð,¨D°°Ð8Ð9Ð9r+   c           
     ó.  — |j                  dd«      }| j                  j                  |«       |› d�}|dk(  rdnd}|rt|d   }	|d   }
|d	   }g }t        t	        |«      «      D ]*  }||
k(  r|j                  |	«       Œ|j                  d
«       Œ, dj                  |«      }|› d|› d|› d|› d|› d�
S |d   }	|d   }|d   }t        j                  j                  |«      }|�t	        |j                  «       «      nd}t	        |«      t	        |«      z   }t	        | j                  «      }|dz
  }||k(  xr ||k(  }|rv|D ��cg c]  \  }}|‘Œ	 }}}t	        |«      }t	        |«      }|dkD  r|dkD  rd|z  }d|z  }|	› d|› d|› d�}n|	}|j                  |«       |j                  d„ |D «       «       n8|D �cg c]  }d‘Œ }}|j                  |	«       |j                  d„ |D «       «       dj                  |«      }|› d|› d|› d|› d|› d�
S c c}}w c c}w )zBBuild store expression for scatter operations (indirect indexing).Úis_point_scatterFrÈ  rÄ  r  rÅ  rÂ  rÁ  Úoutput_shapeÚ0r    rÆ  z	[...].at[z].ré   r!   Údims_beforeÚ
dims_afterr   r   úNone, ú, Nonerj  r  rk  c              3  ó&   K  — | ]	  \  }}|–— Œ y ­wrS   r?   )r@  r_  r
  s      r)   rA  z9PallasKernel._build_scatter_store_expr.<locals>.<genexpr>+  s   è ø€ ÒI©N¨H°dœxÑIùs   ‚ri  c              3  ó    K  — | ]  }d –— Œ y­w)ri  Nr?   )r@  rE  s     r)   rA  z9PallasKernel._build_scatter_store_expr.<locals>.<genexpr>0  s   è ø€ Ò7 qœsÑ7ùs   ‚)r»  ró  r  rK  r  rý  rV  r   r·   r¢  r£  r/  r�  )r$   r¾  r¿  Úscatter_infor£  rÉ  rÎ  rË  rÊ  rÂ  rÁ  rÏ  Úindex_partsrY  Úindex_tuplerÑ  rÒ  rl  Úoutput_ndimÚnum_iter_vars_in_storeÚtotal_kernel_iter_varsÚremaining_dimsÚis_element_wiser_  r
  Ú	n_leadingÚ
n_trailingÚleading_onesÚtrailing_nonesÚindirect_reshapedrE  s                                  r)   Ú_build_scatter_store_exprz&PallasKernel._build_scatter_store_exprç
  s�  € ð (×+Ñ+Ð,>ÀÓFÐð 	×Ñ×"Ñ" 3Ô'Ø˜˜V�nˆð # lÒ2‘U¸ˆ
áà'¨Ñ7ˆLØ'¨Ñ7ˆLØ'¨Ñ7ˆLð ˆKÜœS Ó.Ó/ò ,�Ø˜,Ò&Ø×&Ñ& |Õ4à×&Ñ& sÕ+ð	,ð Ÿ)™) 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ÐFˆ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Ð4ˆKÐ4Ø×Ñ˜|Ô,Ø×ÑÑ7¨JÔ7Ô7à—i‘i Ó,ˆàˆe�8˜K˜=¨	°+°¸bÀÀÈAÈeÈWÐTUÐVð	
ùó+ Gùò  5s   ÅHÆ6	Hc                ó  — | j                   j                  |«      }t        j                  j	                  |«      }|| j
                  |<   | j                  |«      }| j                  |||«      }| j                  ||«      }|�"|| j                  |<   | j                  ||«      }n:| j                  |||«      }| j                  |||«      }| j                  ||||«      }|j                  s4|j                  dk(  r%| j!                  ||«      }| j#                  |||«      }| j$                  j'                  | j(                  ||¬«      }| j                   j                  |«      }	|	| j*                  t-        |«      <   || j.                  |	<   |S )Nr  ©r»   )r#   r¢  r   r·   r  rò  rÝ  rM  rg  r   rp  rO  rT  r‰  rÉ  rÈ  r�  rœ  rÎ   rÏ   rÐ   r  r:   r  )
r$   r£  r  rl  r»   rE  rm  rV  Úcse_varÚ	buf_params
             r)   ÚloadzPallasKernel.load7  sƒ  € à�i‰i�o‰o˜dÓ#ˆÜ—‘×!Ñ! $Ó'ˆð ',ˆ×Ñ˜dÑ#ð ×'Ñ'¨Ó.ˆð ×/Ñ/°°e¸XÓFˆð
 ×/Ñ/°°tÓ<ˆØÐØ/5ˆD×&Ñ& tÑ,Ø×/Ñ/°°VÓ<‰Ið ×:Ñ:¸4ÀÈÓQˆHð ×/Ñ/°°e¸XÓFˆHð ×-Ñ-¨c°4¸ÀÓIˆIð ×%Ò%¨(×*<Ñ*<ÀÒ*EØ×?Ñ?ÀÀiÓPˆIà×7Ñ7¸¸eÀYÓOˆIà—(‘(×#Ñ#Ø�L‰LØØð $ó 
ˆð —I‘I—O‘O DÓ)ˆ	Ø+4ˆ×Ñœ3˜w›<Ñ(Ø/3ˆ×!Ñ! )Ñ,Øˆr+   c                ó	  ‡ ‡‡$‡%— ‰ j                  ‰«      }‰ j                  j                  |«       t        |«      dk(  r‰ j	                  ‰«      S ˆˆ fd„}t        ||d¬«      }|D �cg c]
  } ||«      ‘Œ }}‰ j	                  ‰ j                  ‰«      «      }‰ j                  ‰«      }|D �	cg c]  }	t        |	«      ‘Œ }
}	|D �ci c]  }t        |«       ||«      “Œ }}t        |«      dk(  r¬t        |
«      dk(  rž|d   }t        |«      }|‰ j                  v xr ‰ j                  |   j                  }|rc|‰ j                  v rS‰ j                  |   }|j                  }‰ j                  |«      }d‰ j	                  |«      › d�}|j                  ||«      }|S d}t        |
«      dkD  r!‰ j                  «       |z
  }t        |«      dk(  }|rðdt        |«      z   }|
D ]*  }d	t        |«      z  }|› d
|› d�}|j                  ||«      }Œ, t        |«      D ]£  \  }}t        |«      }|‰ j                  v sŒ ‰ j                  |   }|j                  }‰ j                  |«      }dg|z  }‰ j	                  |«      ||dz   <   dj                  |«      }d‰ j	                  |«      › d|› d�}|j                  ||«      }Œ¥ |S g }|D ]  }|j!                   ||«      d|f«       Œ |D ]  }	|j!                   ||	«      d|	f«       Œ |j#                  d„ d¬«       t        |«      D ]Ç  \  }}t        |«      }|‰ j                  v sŒ ‰ j                  |   }|j                  }‰ j                  |«      } ||«      Š%d‰ j	                  |«      › d�}t%        ˆ%fd„|D «       «      }t%        ˆ%fd„|j'                  «       D «       «      }||z   }|dkD  rd|z  } |› d| › d�}|j                  ||«      }ŒÉ |
D ]‰  }||   Š$t%        ˆ$fd„|D «       «      }!t%        ˆ$fd„|D «       «      }|!dkD  r|dkD  rd|!z  }"d|z  }#|› d|"› d|#› d�}n(|!dkD  rd|!z  }"|› d|"› d�}n|dkD  rd|z  }#|› d|#› d�}n|}|j                  ||«      }Œ‹ |S c c}w c c}	w c c}w )a¡  
        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                ó>   •— ‰j                  ‰| t        d«      ¬«      S )NÚinf)Údefault)r¥  r¨   )rÔ   r  r$   s    €€r)   Ú_coeffz3PallasKernel._handle_mixed_indexing.<locals>._coeff…  s   ø€ Ø×.Ñ.¨u°cÄ5ÈÃ<Ð.ÓPÐPr+   T©ÚkeyrÏ  r   újnp.arange(r!   Fz, 1z.reshape(-1Ú1r    z
).reshape(r  r„  c                ó   — | d   S r>  r?   rZ   s    r)   r„  z5PallasKernel._handle_mixed_indexing.<locals>.<lambda>ê  s
   € ¨!¨A©$€ r+   c              3  ó.   •K  — | ]  }|‰k  sŒ	d –— Œ y­wrD  r?   ©r@  r&   Ú	var_coeffs     €r)   rA  z6PallasKernel._handle_mixed_indexing.<locals>.<genexpr>  s   øè ø€ Ò%N¨AÀÀIÃ¤aÑ%Nùó   ƒ
Žc              3  ó.   •K  — | ]  }|‰k  sŒ	d –— Œ y­w)r   Nr?   rõ  s     €r)   rA  z6PallasKernel._handle_mixed_indexing.<locals>.<genexpr>  s   øè ø€ ò *Ø¸1¸y»=”Añ*ùr÷  rÔ  z[:rk  c              3  ó.   •K  — | ]  }|‰kD  sŒ	d –— Œ y­wrD  r?   ©r@  r&   Úindirect_coeffs     €r)   rA  z6PallasKernel._handle_mixed_indexing.<locals>.<genexpr>  s   øè ø€ ÒI !°a¸.Ó6HœAÑIùr÷  c              3  ó.   •K  — | ]  }|‰k  sŒ	d –— Œ y­wrD  r?   rú  s     €r)   rA  z6PallasKernel._handle_mixed_indexing.<locals>.<genexpr>  s   øè ø€ ÒJ 1°q¸>Ó7IœQÑJùr÷  rÓ  rj  r  z...]z[...)rÊ   rÈ   rÉ   r  rÍ   r  rÌ   r9  r:   r/  r   rO  Úreplacer+  rw  rV  rý  ÚsortrD  rå  )&r$   r  Úused_iter_vars_setrî  rÈ   rÔ   Úiter_coeffsrÈ  Úindirect_var_symsrš  rk  r8  Ú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Ø  z#PallasKernel._handle_mixed_indexingg  s[  û€ ð& "×5Ñ5°eÓ<Ðð 	×Ñ×"Ñ"Ð#5Ô6äÐ!Ó" aÒ'Ø—:‘:˜eÓ$Ð$õ
	Qô  Ð 2¸ÈÔMˆØ.<Ö= s‘v˜c•{Ð=ˆÐ=ð —J‘J˜t×3Ñ3°EÓ:Ó;ˆ	Ø ×3Ñ3°EÓ:ÐØ->Ö? cœ˜S�Ð?ˆÐ?ð 7HÖH°œ3˜q›6¡6¨!£9Ñ,ÐHˆÐ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ð !.ò J�Ø %¬¨NÓ(;Ñ ;�Ø". ¨{¸=¸/ÈÐK�Ø%×-Ñ-¨l¸LÓI‘	ðJô $ NÓ3ò I‘��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ð#Ið& Ðð
 ˆØ!ò 	>ˆCØ×!Ñ!¡6¨#£;°¸Ð"<Õ=ð	>à$ò 	BˆCØ×!Ñ!¡6¨#£;°
¸CÐ"@ÕAð	Bà×Ñ¡¸ÐÔ=ô   Ó/ò 	E‰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	Eð: *ò 	Fˆ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ð+	Fð. Ðùò} >ùò
 @ùò Is   Á(Q;Â.R ÃRc                óF  — |�|dk7  rt        d|› d�«      ‚| j                  j                  |«      }| j                  j	                  |«       t
        j                  j                  |«      }|d uxr t        |j                  «       «      dk(  }|rd|› d�|› d|› d|› d	�g}nÁ| j                  �| j                  ||d
«      }n¡| j                  ||«      }	|	�A| j                  j                  | j                  |«      «       | j!                  |||	||«      g}nL| j#                  |«      }
| j%                  |||
«      }
| j'                  ||
«      }
| j)                  |||||
|«      }|D ]:  }| j*                  j-                  |«       | j.                  j1                  ||f«       Œ< y )NrÄ  zpallas store mode 'z' not supportedr   r¼  r!   r½  z2.shape, _val) if _val.ndim == 0 else _val.reshape(rÇ  F)rV   r#   ÚoutputÚstore_buffer_namesr  r   r·   r¢  r  r£  rö  rÂ  Ú_detect_scatter_patternrÈ   rÉ   rÊ   rä  rÝ  rO  r§  rÌ  Ústoresrx  rñ  rý  )r$   r£  r  r¿  rÉ  r¾  rl  Ú	is_scalarÚstore_linesr×  rE  r`  s               r)   Ústorez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ð ×*Ñ*Ð6Ø"×?Ñ?ÀÀUÈEÓR‘ð  $×;Ñ;¸EÀ4ÓH�àÐ+à×'Ñ'×.Ñ.¨t×/GÑ/GÈÓ/NÔOà×6Ñ6Ø ¨°d¸Dóð#‘Kð  $×3Ñ3°EÓ:�Hð  $×BÑBØ˜e Xó �Hð
  $×9Ñ9¸%ÀÓJ�Hð #'×"8Ñ"8Ø˜T 5¨%°¸4ó#�Kð  ò 	7ˆDØ�K‰K×!Ñ! $Ô'à×"Ñ"×)Ñ)¨3°¨+Õ6ñ	7r+   c                ó    — | j                  |«      }|dk(  rt        j                  | |«      }	 t        |«      S # t        t
        f$ r |cY S w xY w)z=Get integer coefficient of a variable in an index expression.r   )rW  r  ÚdiffrÀ  r"  r#  )r  rÔ   rí  rW  s       r)   r¥  z#PallasKernel._get_index_coefficientf  sQ   € ð
 —‘˜CÓ ˆØ�AŠ:Ü—J‘J˜u cÓ*ˆEð	Ü�u“:ÐøÜœ:Ð&ò 	ØŠNð	ús   ®
9 ¹AÁAc                ó  — | j                  |«      }t        |«      dk7  ry|d   }t        |«      }t        | j	                  ||«      «      }|dk(  ry| j                  |«      s| j                  |||«      S | j                  |||«      S )zDDetect scatter operation pattern. Returns scatter info dict or None.r   Nr   )r9  r  r:   rÀ  r¥  r4  Ú_detect_point_scatterÚ_detect_iter_scatter)r$   r  Úoutput_nameÚindirect_symsÚindirect_symrÂ  rû  s          r)   r  z$PallasKernel._detect_scatter_patterns  s–   € ð ×/Ñ/°Ó6ˆÜˆ}Ó Ò"Øà$ QÑ'ˆÜ˜<Ó(ˆÜ! $×"=Ñ"=¸eÀ\Ó"RÓSˆØ˜QÒØð ×'Ñ'¨Ô.Ø×-Ñ-¨k¸<ÈÓXÐXð ×(Ñ(¨°¸nÓMÐMr+   c                ój  — |sy	 t         j                  j                  |«      }|j                  «       D �cg c]  }t	        |«      ‘Œ }}t        |«      dk  ryd}t        |«      dz
  }t        t        |«      dz
  dd«      D ]  }	||k(  r|	} n
|||	   z  }Œ ||g g d|dœS c c}w # t
        $ r Y yw xY w)z&Detect single-element scatter pattern.Nr   r   r·  T©rÂ  rÁ  rÑ  rÒ  rÎ  rÏ  )r   r·   r¢  r£  rÀ  Ú	Exceptionr  rK  )
r$   r!  rÂ  rû  rl  r8  rÏ  Ú
cumulativerÁ  rY  s
             r)   r  z"PallasKernel._detect_point_scatterˆ  sÞ   € ñ Øð	Ü—'‘'×$Ñ$ [Ó1ˆCØ,/¯L©L«NÖ; qœC �FÐ;ˆLÐ;ô ˆ|Ó˜qÒ Øð ˆ
Ü˜<Ó(¨1Ñ,ˆÜœ˜\Ó*¨QÑ.°°BÓ7ò 	,ˆCØ Ò+Ø"�ÙØ˜, sÑ+Ñ+‰Jð		,ð )Ø(ØØØ $Ø(ñ
ð 	
ùò! <øÜò 	Ùð	ús"   …1B& ¶B!ÁB& Â!B& Â&	B2Â1B2c                ó¸  ‡— | j                  |«      }g }|D ]{  }t        | j                  ||«      «      }|dkD  sŒ$|| j                  v sŒ3| j	                  | j                  |   j
                  «      }|€ y|j                  t        |«      ||f«       Œ} |j                  ‰|df«       |j                  d„ d¬«       t        ˆfd„t        |«      D «       d«      }	|	€yd}
t        ||	dz   d «      D ]  \  }}}||
k7  r y|
|z  }
Œ ||
k7  ry‰|	|d|	 D ���cg c]
  \  }}}||f‘Œ c}}}||	dz   d D ���cg c]
  \  }}}||f‘Œ c}}}d	dd
œS c c}}}w c c}}}w )z0Detect scatter pattern with iteration variables.r   Nr·  c                ó   — | d   S )Nr   r?   rZ   s    r)   r„  z3PallasKernel._detect_iter_scatter.<locals>.<lambda>º  s
   €  A a¡D€ r+   Trï  c              3  ó<   •K  — | ]  \  }\  }}}|‰k(  sŒ|–— Œ y ­wrS   r?   )r@  rÀ  r£  rE  rÂ  s       €r)   rA  z4PallasKernel._detect_iter_scatter.<locals>.<genexpr>¾  s"   øè ø€ ÒR‘?�1‘l�t˜Q ¸TÀ\Ó=QŒQÑRùs   ƒ•r   Fr%  )rÊ   rÀ  r¥  r/  rG  rO  rý  r:   rþ  r  rw  rŽ  )r$   r  rÂ  rû  rÈ   Úall_varsrÔ   rW  rO  Úindirect_posÚexpectedrE  rO  r4  s     `           r)   r   z!PallasKernel._detect_iter_scatter©  sœ  ø€ ð ×1Ñ1°%Ó8ˆð 02ˆØ!ò 	;ˆCÜ˜×3Ñ3°E¸3Ó?Ó@ˆEØ�q‹y˜S D×$9Ñ$9Ò9ØŸ™¨×(=Ñ(=¸cÑ(B×(IÑ(IÓJ�Ø�>ÙØ—‘¤ S£¨5°&Ð 9Õ:ð	;ð 	�‰˜ ~°rÐ:Ô;Ø�‰™.°$ˆÔ7ô ÛR¤i°Ó&9ÔRØó
ˆð ÐØð ˆÜ (¨°,ÀÑ2BÐ2DÐ)EÓ Fò 	ÑˆAˆu�fØ˜Ò ÙØ˜Ñ‰Hð	ð ˜XÒ%Øð )Ø(Ø2:¸=¸LÐ2I×JÐJ¡w q¨!¨Q˜Q šFÔJØ19¸,ÈÑ:JÐ:LÐ1M×NÐN¡g a¨¨A˜A˜qš6ÔNØ %Ø ñ
ð 	
ùô KùÜNs   ÄEÄ6Ec           	     ó†  ‡ — ‰ j                   sJ ‚|dk(  r‰ j                  ||«      S t        |t        «      rt	        d«      ‚|||f}|‰ j
                  j                  v r‰ j
                  j                  |   S ddddddd	d
œ}t        g d¢«      }t        ˆ fd„|D «       «      }‰ j                  |«      }	‰ j                  «       }
t        d„ ‰ j                  j                  «       D «       «      }|xr |	duxr |	dkD  xr	 |
xr |dkD  }|xr |dkD  xr |	du }|dk(  r8|r/‰ j                  «       }t        |«      dk(  r|d   n|}d|› d|› d�}n³d|› d�}n¬|dv r/||   }|r‰ j                  «       }|› d|› d|d   › d�}n‚|› d|› d�}ny||v rL||   }|r1‰ j                  «       }t        |«      dk(  r|d   n|}|› d|› d|› d�}n=|r	|› d|› d�}n2|› d|› d�}n)t	        d|› dt!        |j#                  «       «      › d�«      ‚‰ j
                  j%                  ‰ j&                  ||¬«      }|‰ j
                  j                  |<   |S )aJ  
        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)rD  r¡  rg  ÚminrH  ÚargmaxÚargmin)r[   r>  Úzc              3  ó:   •K  — | ]  }|‰j                   v –— Œ y ­wrS   )r  )r@  r'   r$   s     €r)   rA  z)PallasKernel.reduction.<locals>.<genexpr>  s   øè ø€ ÒI°˜A §¡Ô,ÑIùrô  c              3  ó:   K  — | ]  }|j                   sŒd –— Œ y­wrD  ©r   )r@  rÏ  s     r)   rA  z)PallasKernel.reduction.<locals>.<genexpr>  s   è ø€ ò 
Ø¸U×=OÓ=OŒAñ
ùó   ‚”Nr   r   Úxor_sumzjnp.bitwise_xor.reduce(z, axis=r!   )r1  r2  ré   r·  z, keepdims=True)z	, axis=0)zReduction type 'z8' not yet supported in Pallas backend. Supported types: z	, xor_sumræ  )Úinside_reductionÚwelford_reduce_fallbackr§   r¾  rV   rÎ   Úreduction_cacher   rH  r  rû  rD  r/  rå  r  r  rI  r0  rÏ   rÐ   )r$   r»   r¼   Úreduction_typer¿  Ú	cache_keyÚreduction_opsÚpointwise_prefixesÚhas_pointwiseÚpointwise_numelr#  Ún_reduction_dimsÚis_partial_reductionÚis_symbolic_partialr  Ú	axis_exprÚreduction_exprÚreduction_opr1   s   `                  r)   Ú	reductionzPallasKernel.reductionÖ  sü  ø€ ð" ×$Ò$Ð$Ð$ð Ð-Ò-Ø×/Ñ/°°uÓ=Ð=ä�eœUÔ#ÜØZóð ð
  °Ð6ˆ	Ø˜Ÿ™×0Ñ0Ñ0Ø—8‘8×+Ñ+¨IÑ6Ð6ð ØØØØØ"Ø"ñ
ˆô (ªÓ8ÐÜÓIÐ6HÔIÓIˆØ&*×&@Ñ&@ÐASÓ&TˆØ&*×&CÑ&CÓ&EˆÜñ 
Ø×0Ñ0×7Ñ7Ó9ô
ó 
Ðð
 ò %Ø tÐ+ò%à !Ñ#ò%ð  ò%ð ! 1Ñ$ð 	ð ÒNÐ.°Ñ2ÒN°È$Ð7Nð 	ð ˜YÒ&Ù#Ø×/Ñ/Ó1�Ü'*¨4£y°A¢~˜D šG¸4�	Ø#:¸5¸'ÀÈÈÐSTÐ!U‘à#:¸5¸'ÀÐ!C‘ØÐ3Ñ3Ø(¨Ñ8ˆLÙ#à×/Ñ/Ó1�Ø$0 >°°5°'¸ÀÀbÁÀ
È!Ð!L‘à$0 >°°5°'¸Ð!;‘Ø˜}Ñ,Ø(¨Ñ8ˆLÙ#Ø×/Ñ/Ó1�Ü'*¨4£y°A¢~˜D šG¸4�	à#�n A e W¨G°I°;Ð>NÐOñ ñ %ð %1 >°°5°'¸Ð!C‘à$0 >°°5°'¸Ð!;‘äØ" >Ð"2ð 3$Ü$(¨×);Ñ);Ó)=Ó$>Ð#?¸yðJóð ð —‘×"Ñ"Ø�L‰LØØð #ó 
ˆð /5ˆ�‰× Ñ  Ñ+Øˆr+   c                ó€   — t         j                  j                  | «      }|j                  «       }|j	                  «       S rS   )r   r·   r¢  r‚  r¦  )Úbuffer_namerl  rˆ  s      r)   Ú_buffer_is_contiguousz"PallasKernel._buffer_is_contiguousE  s1   € ä�g‰g× Ñ  Ó-ˆØ—‘Ó!ˆØ×#Ñ#Ó%Ð%r+   c                ó¦	  ‡‡‡ — | j                   ry| j                  ry| j                  ry| j                  r[dj	                  d„ | j
                  j                  D «       «      }| j                  D ]   }|| j                  v rŒt        |«      |v sŒ  y t        | j                  j                  j                  «       «      }d}|D ]�  }| j                  |«      }|€Œ|\  }}}}	}t        |	«      dk\  sŒ.t        |«      dk\  sŒ=|	d   }
|	d   }| j                  |d   «      }| j                  |d   «      }|
€Œr|€Œu|
|k  sŒ{|€Œ~|€Œ�|dkD  sŒ‡|dkD  sŒ�d} n t!        | j"                  «      xs || _        | j'                  «       }|duxr |dkD  }|rt)        d	„ | j*                  D «       «      nd| _        g Š |D ]l  }| j                  |«      }|€ y|\  }}}}}|D �cg c]  }| j                  |«      ‘Œ c}Št/        d
„ ‰D «       «      r yt        ‰«      t        ‰ «      kD  sŒk‰Š Œn ‰ sy| j0                  �t        | j0                  «      Š t        ‰ «      }t        | j                  j2                  «      |z   }d}|D �]
  }| j                  |«      }|€ y|\  }}}}}t        |«      dk(  rŒ0|| j4                  v rt        | j4                  |   «      ŠnY| j0                  �||v rt        | j0                  «      Šn3|D �cg c]  }| j                  |«      ‘Œ c}Št/        d„ ‰D «       «      r yt        ‰«      }||k(  r¹|| j6                  v }d}t9        |«      D ]#  }‰|   ‰ |   k(  s‰|   dk(  s
‰ |   dk(  s|rŒ!d} n |rL|| j"                  v r>| j"                  |   Št        ‰«      |k(  rt;        ˆˆˆ fd„t9        |«      D «       «      s y|r y|s t/        ˆˆ fd„t9        |«      D «       «      s�Œud}�Œy||kD  rSd}t9        ||z
  dz   «      D ]5  }d}t9        |«      D ]  }‰ |   dk(  rŒ‰||z      ‰ |   k7  sŒd} n |sŒ3d} n |s yd}�ŒÑt=        t?        ‰«      t?        ‰ «      «      D ]  \  }}||k7  sŒ|dk7  sŒ|dk7  sŒ  y �Œ |syt@        jB                  jE                  «       jF                  dk(  }|sd| _$        yd| _$        yc c}w c c}w )a   Check if this kernel can use tiling on CPU/TPU.

        Tiling is compatible with reductions, transpositions, and multi-range-tree
        kernels as long as no flatten-based indexing is used (buf[...].flatten()[idx]).
        Flatten indexing requires global flat indices which don't decompose into
        per-tile local indices.

        Reject:
        - GPU (has its own TMA / padding path)
        - Flatten-based indexing
        - Scatter outputs (indirect indexing complicates tile boundaries)
        Fú
c              3  ó2   K  — | ]  }t        |«      –— Œ y ­wrS   ©r:   ©r@  r`  s     r)   rA  z1PallasKernel._can_tile_cpu_tpu.<locals>.<genexpr>d  s   è ø€ Ò$O°4¤S¨§YÑ$OùrB  Nr   r   r   Tc              3  ó:   K  — | ]  }|j                   sŒd –— Œ y­wrD  r6  )r@  r  s     r)   rA  z1PallasKernel._can_tile_cpu_tpu.<locals>.<genexpr>�  s   è ø€ ÒD�d°$×2CÓ2C”ÑDùr7  c              3  ó$   K  — | ]  }|d u –— Œ
 y ­wrS   r?   r?  s     r)   rA  z1PallasKernel._can_tile_cpu_tpu.<locals>.<genexpr>œ  s   è ø€ Ò/ �1˜”9Ñ/ùrš  c              3  ó$   K  — | ]  }|d u –— Œ
 y ­wrS   r?   r?  s     r)   rA  z1PallasKernel._can_tile_cpu_tpu.<locals>.<genexpr>¾  s   è ø€ Ò3 Q�q˜D”yÑ3ùrš  c              3  óh   •K  — | ])  }‰‰|      ‰|   k(  xs ‰‰|      d k(  xs ‰|   d k(  –— Œ+ y­wrD  r?   )r@  rÀ  Úint_sizer¯  Ú	ref_shapes     €€€r)   rA  z1PallasKernel._can_tile_cpu_tpu.<locals>.<genexpr>Ö  sY   øè ø€ ò  ð !"ð % T¨!¡WÑ-°¸1±Ñ=ò 1Ø'¨¨Q©Ñ0°AÑ5ò1à(¨™|¨qÑ0ó1ñ ùs   ƒ/2c              3  óH   •K  — | ]  }‰|   ‰|   k(  xr ‰|   d kD  –— Œ y­wrD  r?   )r@  rÀ  rU  rV  s     €€r)   rA  z1PallasKernel._can_tile_cpu_tpu.<locals>.<genexpr>â  s7   øè ø€ ò %àð ˜Q‘K 9¨Q¡<Ñ/ÒD°I¸a±LÀ1Ñ4DÓDñ%ùrø  r³   é@   )%rï  rÿ  ró  rÈ   rV  rÐ   rZ  rþ  r:   rI  r#   Úoutput_buffersr0  rF  r  rG  rÛ   rô  Útile_has_transposerû  rD  r  Útile_skip_last_nrH  rö  r!  rõ  r   rK  rL  ÚziprŽ  r   r·   r¸   r¹   r÷  )!r$   Úcompute_textrÎ  Úout_bufsÚhas_col_major_outr{  rM   rE  rP  r  r³  r´  Úd0Úd1r#  Úhas_reductionr8  Úref_ndÚall_bufsÚhas_tileableÚbuf_ndÚ
is_stridedÚmismatchrÀ  Úfoundr´  Úokr˜   r™   rÖ  rU  r¯  rV  s!                                 @@@r)   Ú_can_tile_cpu_tpuzPallasKernel._can_tile_cpu_tpuK  sN  ú€ ð �;Š;ØØ×$Ò$ØØ×!Ò!Øð ×ÒØŸ9™9Ñ$O¸4¿<¹<×;NÑ;NÔ$OÓOˆLØ×.Ñ.ò !�Ø˜d×:Ñ:Ñ:ØÜ�w“< <Ò/Ù ð	!ô ˜Ÿ	™	×0Ñ0×5Ñ5Ó7Ó8ˆð "ÐØ ò 	ˆHØ×(Ñ(¨Ó2ˆDØˆ|ØØ04Ñ-ˆAˆx˜˜N¨AÜ�>Ó" aÓ'¬C°«M¸QÓ,>Ø# AÑ&�Ø# AÑ&�Ø—^‘^ H¨Q¡KÓ0�Ø—^‘^ H¨Q¡KÓ0�à‘NØ™Ø˜R›Ø™Ø™Ø˜Q›Ø˜Q›à(,Ð%Ùð)	ô* #' t×'BÑ'BÓ"CÒ"XÐGXˆÔð
 ×7Ñ7Ó9ˆØ'¨tÐ3ÒK¸È!Ñ8Kˆñ ô ÑD˜d×.Ñ.ÔDÔDàð 	Ôð  "ˆ	Ø ò 		%ˆHØ×(Ñ(¨Ó2ˆDØˆ|ÙØ#'Ñ ˆAˆx˜˜A˜qØ3;Ö<¨a˜Ÿ™ qÕ)Ò<ˆHÜÑ/ hÔ/Ô/ÙÜ�8‹}œs 9›~Ó-Ø$‘	ð		%ñ Øð
 ×&Ñ&Ð2Ü˜T×8Ñ8Ó9ˆIä�Y“ˆä˜Ÿ	™	×/Ñ/Ó0°8Ñ;ˆØˆØ ó R	%ˆHØ×(Ñ(¨Ó2ˆDØˆ|ÙØ#'Ñ ˆAˆx˜˜A˜qÜ�8‹} Ò!Øð ˜4×8Ñ8Ñ8Ü × =Ñ =¸hÑ GÓH‘Ø×,Ñ,Ð8¸XÈÑ=QÜ × ;Ñ ;Ó<‘à7?Ö@°!˜DŸN™N¨1Õ-Ò@�ÜÑ3¨(Ô3Ô3Ù Ü˜“]ˆFà˜Òð &¨×)CÑ)CÐC�
Ø �Ü˜v›ò 	�Aà  ™ y°¡|Ò3Ø# A™;¨!Ò+Ø$ Q™<¨1Ò,Ù%à Ø#�HÙð	ñ  ¨D×,GÑ,GÑ GØ×6Ñ6°xÑ@�Dä˜D›	 VÒ+Üõ  ô &+¨6£]ô	 ô ñ  %ÙÙ ñ ¤ô %ä" 6›]ô%ö "ð $(’Là˜&’ð �Ü˜v¨™°Ñ2Ó3ò 
�AØ�BÜ" 6›]ò "˜Ø$ Q™<¨1Ò,Ø$Ø# A¨¡E™?¨i¸©lÓ:Ø!&˜BÙ!ð"ò Ø $˜Ùð
ñ Ù Ø#’ô  ¤¨Ó 2´H¸YÓ4GÓHò %‘D�A�qØ˜A“v ! q£&¨Q°!«VÚ$ò%ðaR	%ñh Øô —‘×4Ñ4Ó6×;Ñ;¸uÑDˆÙØ)+ˆDÔ&ð ð *.ˆDÔ&àùòi =ùòD As   ÇS	Ë<Sc           
     ó†  — t        «       }| j                  j                  «       \  }}}}|D �cg c]  }|j                  ‘Œ }}|D �cg c]  }|j	                  d«      sŒ|‘Œ }	}|D �cg c]  }|j	                  d«      sŒ|‘Œ }
}t        | j                  j                  j                  «       «      }|D �cg c]	  }||v sŒ|‘Œ }}|
st        d«      ‚| j                  j                  j                  «       D ��ci c]  \  }}t        |t        «      r||“Œ }}}|xs d}t        j                  j                  «       j                   dk(  }|rdnd}i }|	D ]  }d||<   Œ	 |	D �cg c]  }||   sŒ	|› d	�‘Œ }}|D �cg c]  }|j	                  d
«      sŒ|‘Œ }}||z   }||z   }t        |j                  «       D ��cg c]
  \  }}|rŒ	|‘Œ c}}«      }|rt#        t%        t'        |
«      «      «      }n0| j(                  rg }n!t+        |
«      D ��cg c]  \  }}||v r|‘Œ }}}t-        dNi d|“d|“d| j(                  “d|“d|“d|“d|	“d|
“d|“d|“d|“d|“d|“d|“d|“d|“d|“dg “Ž}|| _        | j1                  |«       t        «       }|j3                  «       5  | j5                  ||«       | j6                  j8                  D ]  }|j;                  t        |«      «       Œ 	 ddd«       | j                  j                  «       \  }}}}|D �cg c]  }|j                  ‘Œ }}t        | j                  j                  j                  «       «      }|D �cg c]	  }||v sŒ|‘Œ c}|_        |D �cg c]  }|j	                  d
«      sŒ|‘Œ c}|_        ||j>                  z   |_         ||z   |_!        | jE                  «       | _#        d} | jH                  rd} | jK                  ||«      |_&        tO        | jP                  «      }!|!rŒ| jS                  «        t        «       }"|"j3                  «       5  | j6                  j8                  D ]  }|"j;                  t        |«      «       Œ 	 ddd«       | jU                  |||"«       |jW                  «       S d |› d!d"jY                  |jB                  «      › | › d#�}#|j;                  |#«       |j3                  «       5  | j[                  |||«       ddd«       |j;                  d«       |› d$�}$g }%d%t'        |j<                  «      z   }&t+        |j@                  «      D ]/  \  }}||v s|j	                  d&«      sŒ|%j]                  ||&z   «       Œ1 |%rd'd"jY                  d(„ |%D «       «      z   d)z   }'nd*}'t#        t%        d%t'        |j<                  «      z   «      «      }(d'd"jY                  d+„ |(D «       «      z   d)z   })|j;                  d,|)› d-|'› d.�«       d/d0g|j<                  z   |j@                  z   }*|j;                  d |$› d'd"jY                  |*«      › d#�«       d"jY                  d1„ |jL                  D «       «      }+| j_                  |«      \  },}-d2}.|j3                  «       5  |,r|ja                  |.«       �nŒ|-r;|j;                  d3«       |j3                  «       5  |ja                  |.«       ddd«       |j;                  d4«       | jb                  �|j;                  d5| jb                  › d)�«       |jL                  D ].  \  }/}0|j@                  |/   }|j;                  |› d6|› d7|0› d8�«       Œ0 |j;                  d9«       |j;                  d:«       |j;                  d;«       |j;                  d.«       | jF                  r| je                  |«       �n!| jg                  ||j@                  «       |j@                  D ]O  }| ji                  |«      }1|1r| jj                  jm                  |1«      nd}2|2€Œ6|j;                  |› d6|› d<|2› d.�«       ŒQ |j;                  d=«       |j;                  d>«       |j;                  d;«       |j;                  d.«       |j;                  d?«       |j;                  d@«       |j;                  dAd"jY                  |j@                  «      z   dBz   «       |j;                  d.«       | jH                  rQ| jF                  r#|j;                  dC«       |j;                  dD«       n"|j;                  dE«       |j;                  dF«       g }3|j<                  D ]  }4|3j]                  |4› dG|4› �«       Œ | jH                  r"|3j]                  dH«       |3j]                  dI«       |3rdJ|› dKd"jY                  |3«      › dL�}5n|› dM�}5| jn                  xr | jp                  xr | js                  «       }6|6r| ju                  ||5«       n=| jn                  r| jw                  ||5«       n| jy                  ||5|jL                  |+«       ddd«       | j{                  ||$«       |jW                  «       S c c}w c c}w c c}w c c}w c c}}w c c}w c c}w c c}}w c c}}w # 1 sw Y   �Œ‘xY wc c}w c c}w c c}w # 1 sw Y   �Œ-xY w# 1 sw Y   �Œ¶xY w# 1 sw Y   �ŒéxY w# 1 sw Y   ŒžxY w)OaÍ  
        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)rm  Ú
in_out_ptrz2Pallas backend requires at least one output bufferú<KERNEL_NAME>ÚcpurÙ   rÚ   TrÈ  )rn  Úin_ptrrÔ  rÕ  rÖ  r×  rØ  rÚ  rÛ  rÜ  rÝ  rÞ  rà  rá  râ  rã  rä  rå  rç  ré  NrR  z&, _pallas_tile=None, _pallas_ax2g=Noneúdef ú_kernel(r    ú):Ú_jit_wrapperr   rn  ré   c              3  ó2   K  — | ]  }t        |«      –— Œ y ­wrS   rO  ©r@  r[   s     r)   rA  z.PallasKernel.codegen_kernel.<locals>.<genexpr>¬  s   è ø€ Ò,L¸¬S°¯VÑ,LùrB  ú,)z()c              3  ó2   K  — | ]  }t        |«      –— Œ y ­wrS   rO  rw  s     r)   rA  z.PallasKernel.codegen_kernel.<locals>.<genexpr>°  ó   è ø€ Ò0P¸A´°Q·Ñ0PùrB  ú+@functools.partial(jax.jit, static_argnums=z, donate_argnums=r!   Ú
out_shapesÚ
out_dtypesc              3  ó0   K  — | ]  \  }}|› d |› �–— Œ y­w©ú: Nr?   ©r@  rÀ  Úos      r)   rA  z.PallasKernel.codegen_kernel.<locals>.<genexpr>»  s   è ø€ Ò%P±f°q¸!¨¨¨2¨a¨S¤kÑ%Pùó   ‚)zPresults = tuple(jnp.empty(s, dtype=dt) for s, dt in zip(out_shapes, out_dtypes))z2return results if len(results) > 1 else results[0]z+if any(0 in shape for shape in out_shapes):zI_pallas_out_shapes = tuple(s if len(s) > 0 else (1,) for s in out_shapes)z_pallas_out_shapes = (rr  z.reshape(_pallas_out_shapes[ú])zout_shapes_pallas = tuple(z&    jax.ShapeDtypeStruct(shape, dtype)z;    for shape, dtype in zip(_pallas_out_shapes, out_dtypes)rs  úout_specs_pallas = tuple(z+    pallas_make_block_spec_non_tiled(shape)úin_specs_pallas = tuple(z-    pallas_make_block_spec_non_tiled(i.shape)z    for i in [rk  z_pallas_tile = _tilez_pallas_ax2g = _ax2gz$_pallas_tile = _pallas_out_shapes[0]z_pallas_ax2g = {}ú=z_pallas_tile=_pallas_tilez_pallas_ax2g=_pallas_ax2gúfunctools.partial(ú	_kernel, ú),z_kernel,r?   )>r   r#   Úpython_argdefsr£  r7  r   r  rå  ÚRuntimeErrorrY  rN  r§   r:   r   r·   r¸   r¹   rI  rK  r  rÖ  rw  rÓ  Úaliasable_out_ptrsÚ_codegen_importsÚindentÚ_codegen_iteration_varsrÐ   rZ  rx  rÝ  râ  rã  rä  rk  Útile_cpu_tpurþ  Ú_compute_alias_pairsré  rÛ   r  r~  Ú _codegen_scalar_prefetch_wrapperÚgetvaluerV  Ú_emit_kernel_bodyrý  rê  Ú
writelinesrö  Ú_codegen_tiled_specsr€  rv  rõ  r»  rï  rð  r*  Ú_codegen_jit_wrapper_tmaÚ_codegen_jit_wrapper_legacy_gpuÚ_codegen_jit_wrapper_cpu_tpuÚ_codegen_main_entry)7r$   r£  rÔ  Úarg_defsÚ	call_argsrE  r˜   rÚ  r'   rÛ  rÜ  Úsize_var_namesrÝ  Úouterr  rÞ  rÕ  r×  rØ  rà  rz  rá  râ  rã  rä  Úflagrå  rç  rÌ  rç  Úkernel_bodyr`  Úextra_kernel_paramsÚuse_scalar_prefetchÚkernel_body_spÚkernel_signatureÚjit_wrapper_nameÚdonate_indicesÚbase_offsetÚdonate_literalÚstatic_argnumsÚstatic_argnums_literalÚwrapper_paramsÚalias_map_literalÚhas_zero_dimrè  Úzero_dim_returnÚ	input_idxÚout_idxr{  ÚcshapeÚpartial_argsÚsv_paramÚ
kernel_argÚuse_tmas7                                                          r)   Úcodegen_kernelzPallasKernel.codegen_kernel  s
  € ô Óˆð %)§I¡I×$<Ñ$<Ó$>Ñ!ˆ�)˜Q Ø)1Ö2 A˜Ÿ›Ð2ˆÐ2Ø&3ÖO °q·|±|ÀIÕ7Nš1ÐOˆÐOà$ö
Ø¨¯©Ð5NÕ(OŠAð
ˆð 
ô $ D§I¡I×$6Ñ$6×$=Ñ$=Ó$?Ó@ˆØ&3ÖK °q¸NÒ7Jš1ÐKˆÐKÙÜÐSÓTÐTð !%§	¡	× 8Ñ 8× >Ñ >Ó @÷ 
á��uÜ˜%¤Ô%ð �5‰Lð 
Ðñ  
ð Ò-˜oˆÜŸ7™7×>Ñ>Ó@×EÑEÈÑNÐÙ&6™F¸GÐà+-ˆØ$ò 	*ˆEØ%)ˆO˜EÒ"ð	*ð +:ö
Ø!&¸_ÈUÓ=Sˆuˆg�VÒð
ˆð 
ð %ö
Ø¨¯©Ð5MÕ(NŠAð
ˆð 
ð +¨\Ñ9ÐØ)¨MÑ9ÐÜ&Ø$3×$9Ñ$9Ó$;×H‘j�d˜DÂ4ŠTÓHó
Ðñ Ü"&¤u¬S°Ó-?Ó'@Ó"AÑØ�[Š[Ø"$Ñô "+¨=Ó!9÷#á�C˜ØÐ,Ñ,ò ð#Ðñ #ô ò 
Ùð
á#ð
ð —;’;ð
ñ .ð	
ñ
 0ð
ñ (ð
ñ ,ð
ñ (ð
ñ ,ð
ñ "6ð
ñ ,ð
ñ &ð
ñ &ð
ñ !4ð
ñ  2ð
ñ  0ð!
ñ" !4ð#
ñ$ ð%
ˆð( #2ˆÔà×Ñ˜cÔ"ä$Ó&ˆØ×ÑÓ!ñ 	1Ø×(Ñ(¨°cÔ:àŸ™×+Ñ+ò 1�Ø×%Ñ%¤c¨$£iÕ0ñ1÷	1ð %)§I¡I×$<Ñ$<Ó$>Ñ!ˆ�)˜Q Ø)1Ö2 A˜Ÿ›Ð2ˆÐ2Ü# D§I¡I×$6Ñ$6×$=Ñ$=Ó$?Ó@ˆØ*7ÖO Q¸1ÀÒ;NšqÒOˆÔà$ö
Ø¨¯©Ð5MÕ(NŠAò
ˆÔð #/°×1AÑ1AÑ"AˆÔØ!-°Ñ!=ˆÔð !×2Ñ2Ó4ˆÔà ÐØ×'Ò'Ø"JÐà×3Ñ3°C¸ÓIˆŒä" 4×#7Ñ#7Ó8ÐáØ×.Ñ.Ô0Ü+Ó-ˆNØ×&Ñ&Ó(ñ 8Ø ŸL™L×/Ñ/ò 8�DØ"×,Ñ,¬S°«YÕ7ñ8÷8ð ×1Ñ1ØØØôð
 —=‘=“?Ð"ð " + ¨h°t·y±yÀ×AWÑAWÓ7XÐ6YÐZmÐYnÐnpÐqÐØ�‰Ð'Ô(à�[‰[‹]ñ 	;Ø×"Ñ" 4¨°cÔ:÷	;ð 	�‰�rÔØ)˜]¨,Ð7ÐØˆØœ#˜c×1Ñ1Ó2Ñ2ˆÜ" 3×#:Ñ#:Ó;ò 	9‰IˆC�Ø˜Ñ$¨¯©¸Õ)FØ×%Ñ% c¨KÑ&7Õ8ð	9ñ Ø  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à ŸI™IÑ%PÀÇÁÔ%PÓPÐà(,×(CÑ(CÀCÓ(HÑ%ˆ�oð
ˆð �[‰[‹]ñ Z	ÙØ—‘ Ö0á"Ø—N‘NÐ#PÔQØŸ™›ñ 9ØŸ™¨Ô8÷9ð —‘ðEôð ×.Ñ.Ð:Ø—N‘NØ0°×1LÑ1LÐ0MÈRÐPôð +.¯/©/ò Ñ&�I˜wØ×3Ñ3°IÑ>�EØ—N‘NØ ˜'  U GÐ+GÈÀyÐPRÐSõðð
 —‘Ð;Ô<Ø—‘ÐGÔHØ—‘ØQôð —‘˜sÔ#Ø×$Ò$Ø×-Ñ-¨cÖ2à×2Ñ2°4¸×9PÑ9PÔQØ!$×!8Ñ!8ò S˜Ø#'×#:Ñ#:¸5Ó#A˜ñ  (ð !×9Ñ9×=Ñ=¸hÔGà!%ð ð
 "Ñ-Ø ŸN™N¨e¨W°C¸°w¸iÈÀxÈqÐ+QÕRðSð —N‘NÐ#>Ô?Ø—N‘NÐ#PÔQØ—N‘NØUôð —N‘N 3Ô'Ø—N‘NÐ#=Ô>Ø—N‘NÐ#RÔSØ—N‘NØ(¨4¯9©9°S×5LÑ5LÓ+MÑMÐPSÑSôð —N‘N 3Ô'à×/Ò/Ø×(Ò(ØŸ™Ð'=Ô>ØŸ™Ð'=Õ>àŸ™Ð'MÔNØŸ™Ð':Ô;ð  "�Ø #× 3Ñ 3ò B�HØ ×'Ñ'¨8¨*°A°h°ZÐ(@ÕAðBð ×/Ò/Ø ×'Ñ'Ð(CÔDØ ×'Ñ'Ð(CÔDáØ#5°k°]À)ÈDÏIÉIÐVbÓLcÐKdÐdfÐ!g‘Jà$/ =°Ð!9�Jð —K‘Kò 5Ø×.Ñ.ò5à×2Ñ2Ó4ð ñ
 Ø×1Ñ1°#°zÕBØ—[’[Ø×8Ñ8¸¸jÕIà×5Ñ5Ø˜Z¨¯©Ð:Kô÷qZ	ðx 	× Ñ  Ð&6Ô7Ø�}‰}‹Ðùò} 3ùÚOùò
ùò
 Lùó 
ùò
ùò
ùó Iùó#÷>	1ñ 	1üò 3ùâOùò
÷*8ñ 8ú÷	;ñ 	;ú÷T9ñ 9ú÷Z	ð Z	úsÈ   ®iÁi	Ái	Á(iÁ?iÂ6	iÃ iÃ:iÅ#
iÅ.iÅ;i#Æi#Æ:
i(
Çi(
Èi.ÊAi4ÌjÍ	jÍ jÍ0jÎjÐ.6jÓjÙ89j7Ú1j*ÛEj7àHj7é4i>êjêj'ê*j4	ê/j7ê7k c                óP  — g }t        | j                  «      D ]‹  \  }}|j                  d«      rG|j                  |d«      sŒ*|› d�}| j                  j                  |«      }|j                  ||f«       Œ^| j                  j                  |«      }|j                  ||f«       Œ� |S )Nrm  FrÈ  )rw  rÜ  r7  r»  rã  r  rý  )rç  rà  ré  r±  r£  Ú
alias_namer°  s          r)   r’  z!PallasKernel._compute_alias_pairs$  s¨   € ð .0ˆÜ& s×'8Ñ'8Ó9ò 	9‰MˆG�TØ�‰˜yÔ)Ø"×&Ñ& t¨UÕ3Ø$( 6¨ �JØ #× 7Ñ 7× =Ñ =¸jÓ I�IØ×&Ñ&¨	°7Ð';Õ<à×3Ñ3×9Ñ9¸$Ó?�	Ø×"Ñ" I¨wÐ#7Õ8ð	9ð Ðr+   c                ó$  — |j                   D ]M  }t        |t        «      r |j                  |j	                  «       «       Œ3|j                   j                  |«       ŒO | j                  D ]%  \  }}||j                  v sŒ|j                  |«       Œ' y)z:Emit the kernel body lines and store operations into code.N)rZ  r§   r:   rx  rÙ  rý  rñ  rä  )r$   rÔ  r¡  rç  r`  rm  ry  s          r)   r•  zPallasKernel._emit_kernel_body4  s}   € ð  ×&Ñ&ò 	)ˆDÜ˜$¤Ô$Ø—‘˜tŸ{™{›}Õ-à—‘×"Ñ" 4Õ(ð		)ð
 $(×#9Ñ#9ò 	+ÑˆG�ZØ˜#×0Ñ0Ò0Ø—‘˜zÕ*ñ	+r+   c                ó0  — | j                   €J ‚| j                   }|j                  }t        |j                  «      }|j                  D �cg c]&  }||j
                  k7  r||j                  k7  r||vr|‘Œ( }}|j
                  g|j                  gz   |z   t        |j                  «      z   |j                  z   }	|j                  d|› ddj                  |	«      › d�«       |j                  «       5  | j                  |||«       ddd«       |j                  d«       |› d�}
dd	g|j                  z   |j                  z   }t        t        d
t        |j                  «      z   «      «      }ddj                  d„ |D «       «      z   dz   }|j                  d|› d�«       |j                  d|
› ddj                  |«      › d�«       |j                  «       5  |j                  }|j
                  }|j                   }t        |j"                  «      }|j                  d«       t        |«      D ]   }||k7  sŒ	|j                  d|› d|› d�«       Œ" |j                  d|› d�«       |dk(  r|j                  d|› d|› d�«       nA|gt        |«      D �cg c]
  }||k7  sŒ	|‘Œ c}¢­}|j                  d|› d|› d|› d|› d�	«       g }|D ]¢  }d|› d�}|j                  d|› d|› d �«       |j                  d!|› d"|› d#�«       |j                  d$|› d%�«       |j                  d!|› d"|› d&|› d'�«       |j                  d(«       |j                  d!|› d"|› d)�«       |j%                  |«       Œ¤ g }|j                  D ]0  }d|› d�}|j                  |› d"|› d#�«       |j%                  |«       Œ2 |j                  D �cg c]	  }|› d*|› �‘Œ }}|rd+|› d,dj                  |«      › d�}n|› d-�}|j                  d.«       d/t        |«      z   }|j                  d0«       |j                  «       5  |j                  d1«       |D ]  }|j                  d2|› d3|› d4�«       Œ |j                  D ]  }|j                  d5«       Œ 	 ddd«       |j                  d«       t        |j                  «      }|j                  d6dj                  d7g|z  «      z   dz   «       g }t'        |j                  «      D ]#  \  }}d/|z   |z   } |j%                  | › d8|› �«       Œ% dj                  |«      }!t        |«      D �cg c]  }d9|› d:�‘Œ
 }"}d;dj                  |"«      z   dz   }#|j                  d<«       |j                  «       5  |j                  |› d=�«       |j                  d>|#› d=�«       |j                  d?«       |j                  «       5  |j                  d@«       |j                  dA«       |j                  dB«       |j                  dC«       ddd«       |j                  dD«       |r|j                  dE|!› dF�«       | j(                  s|j                  dG|j*                  › d=�«       ddd«       |gdHgz   |z   |z   }$|j                  dIdj                  |$«      › d�«       |j                  dJ«       ddd«       | j-                  ||
«       yc c}w # 1 sw Y   �ŒÊxY wc c}w c c}w # 1 sw Y   �ŒhxY wc c}w # 1 sw Y   ŒïxY w# 1 sw Y   Œ xY w# 1 sw Y   ŒbxY w)Kz=Emit kernel, JIT wrapper, and main entry for scalar prefetch.Nrr  rs  r    rt  rR  ru  r|  r}  r   ré   c              3  ó2   K  — | ]  }t        |«      –— Œ y ­wrS   rO  rw  s     r)   rA  z@PallasKernel._codegen_scalar_prefetch_wrapper.<locals>.<genexpr>n  rz  rB  rx  r{  r!   z_D = 1z
_D = _D * z.shape[rk  z_seq = z	.shape[0]r   z_table_3d = rs  z.shape[0], 1, _D)z
.transposez	], 1, _D)rE  Ú_3dzif z.ndim == 2 and z.shape[0] == _seq:ú    rr  z.reshape(_seq, 1, _D)zelif z.ndim == 1:z.reshape(1, 1, z
.shape[0])úelse:z.reshape(1, 1, -1)r‡  rˆ  r‰  Ú_kernelzR_ROW_SPEC = pl.BlockSpec((1, 1, _D), lambda i, _: (i, jnp.int32(0), jnp.int32(0)))r   z_in_specs = [zPpl.BlockSpec((1, 1, _D), lambda gi, idx: (idx[gi], jnp.int32(0), jnp.int32(0))),z_ROW_SPEC if z$.shape[0] == _seq else pl.BlockSpec(zA.shape, lambda i, _: (jnp.int32(0), jnp.int32(0), jnp.int32(0))),z
_ROW_SPEC,z_out_specs = [Ú	_ROW_SPECr€  z/jax.ShapeDtypeStruct((_seq, 1, _D), out_dtypes[r„  rj  ú_result = pl.pallas_call(ú,z
out_shape=z'grid_spec=pltpu.PrefetchScalarGridSpec(znum_scalar_prefetch=1,zgrid=(_seq,),zin_specs=_in_specs,zout_specs=_out_specs,rŠ  zinput_output_aliases={ ú },z
interpret=Ú	_table_3dú)(z?return tuple(r.reshape(s) for r, s in zip(_result, out_shapes)))r  rÔ  r   rá  rã  rÃ  r¼  rI  rÜ  rx  rV  r�  r•  rÝ  rK  r  rÁ  r¿  rý  rw  rÖ  rØ  r›  )%r$   rç  rÕ  r¡  r„  rÔ  Ú	alias_setr'   Úother_input_paramsÚprefetch_kernel_paramsr¦  r¬  rª  r«  ÚtableÚindicesÚind_dimrS  rÀ  rZ  r¯  Úpallas_call_other_argsÚp3dÚpallas_call_alias_argsÚsvr³  Ú
kernel_refÚnum_non_alias_in_specsrE  Únum_outputsÚalias_map_partsr±  Úarg_idxr­  Úout_shape_partsÚout_shape_exprÚall_pallas_argss%                                        r)   r“  z-PallasKernel._codegen_scalar_prefetch_wrapperD  sÄ  € ð ×#Ñ#Ð/Ð/Ð/Ø×'Ñ'ˆØ�x‰xˆä˜s×/Ñ/Ó0ˆ	ð ×,Ñ,ö
àØ�H×*Ñ*Ò*Ø�X×)Ñ)Ò)Ø˜Ñ"ò	 ð
Ðð 
ð ×#Ñ#Ð$Ø×#Ñ#Ð$ñ%à ñ!ô �3×#Ñ#Ó$ñ%ð ×Ññ	 ð 	ð 	�‰Ø�;�-˜x¨¯	©	Ð2HÓ(IÐ'JÈ"ÐMô	
ð �[‰[‹]ñ 	;Ø×"Ñ" 4¨°cÔ:÷	;ð 	�‰�rÔØ)˜]¨,Ð7Ðà˜<Ð(¨3×+>Ñ+>Ñ>À×AXÑAXÑXð 	ô œe A¬¨C×,?Ñ,?Ó(@Ñ$@ÓAÓBˆØ!$ t§y¡yÑ0PÀÔ0PÓ'PÑ!PÐSWÑ!WÐØ�‰Ø9Ð:PÐ9QÐQRÐSô	
ð 	�‰˜Ð.Ð/¨q°·±¸>Ó1JÐ0KÈ2ÐNÔOà�[‰[‹]ñ z	Ø×(Ñ(ˆEØ×,Ñ,ˆGà×+Ñ+ˆGÜ�x×+Ñ+Ó,ˆDØ�N‰N˜8Ô$Ü˜4“[ò D�Ø˜“<Ø—N‘N Z°¨w°g¸a¸SÀÐ#BÕCðDð �N‰N˜W W I¨YÐ7Ô8à˜!Š|Ø—‘ ¨e¨W°I¸e¸WÐDUÐVÕWàÐK¬e°D«kÖ"J¨¸QÀ'»\¢1Ò"JÑK�Ø—‘Ø" 5 '¨°D°6¸Ø�g˜W W I¨Yð8ôð &(Ð"Ø'ò 3�Ø˜!˜˜C�j�Ø—‘  Q C °q°cÐ9KÐLÔMØ—‘  c U¨#¨a¨SÐ0EÐFÔGØ—‘  q c¨Ð5Ô6Ø—‘  c U¨#¨a¨S°À¸sÀ*ÐMÔNØ—‘˜wÔ'Ø—‘  c U¨#¨a¨SÐ0BÐCÔDØ&×-Ñ-¨cÕ2ð3ð &(Ð"Ø×%Ñ%ò 3�Ø˜!˜˜C�j�Ø—‘ #  c¨!¨Ð,AÐBÔCØ&×-Ñ-¨cÕ2ð3ð
 47×3FÑ3FÖG¨R˜r˜d ! B 4šLÐGˆLÐGÙà(¨¨ð 6ØŸ	™	 ,Ó/Ð0°ð3ñ ð
 !,˜}¨GÐ4�
ð �N‰NðAôð
 &'¬Ð-CÓ)DÑ%DÐ"Ø�N‰N˜?Ô+Ø—‘“ñ 1Ø—‘ðOôð 2ò �CØ—N‘NðØ"˜eð $)Ø),¨ð .UðVõðð ×)Ñ)ò 1�AØ—N‘N <Õ0ñ1÷1ð �N‰N˜3Ôä˜c×/Ñ/Ó0ˆKØ�N‰NØ  4§9¡9¨k¨]¸[Ñ-HÓ#IÑIÈCÑOôð !ˆOÜ'¨×(8Ñ(8Ó9ò @‘
�˜ØÐ4Ñ4°wÑ>�Ø×&Ñ&¨'¨°"°W°IÐ'>Õ?ð@ð !%§	¡	¨/Ó :Ðô ˜{Ó+öàð BÀ!ÀÀBÒGðˆOð ð ! 4§9¡9¨_Ó#=Ñ=ÀÑCˆNà�N‰NÐ6Ô7Ø—‘“ñ JØ—‘ * ¨QÐ/Ô0Ø—‘ ¨NÐ+;¸1Ð=Ô>Ø—‘ÐHÔIØ—[‘[“]ñ <Ø—N‘NÐ#;Ô<Ø—N‘N ?Ô3Ø—N‘NÐ#8Ô9Ø—N‘NÐ#:Ô;÷	<ð
 —‘˜tÔ$Ù"Ø—N‘NÐ%=Ð>OÐ=PÐPTÐ#UÔVØ—{’{Ø—N‘N Z°×0EÑ0EÐ/FÀaÐ#HÔI÷Jð  �	Ø�-ñ à(ñ)ð )ñ)ð ð �N‰N˜R §	¡	¨/Ó :Ð;¸1Ð=Ô>à�N‰NØQô÷qz	ðx 	× Ñ  Ð&6Õ7ùòA
÷(	;ñ 	;üò> #Kùò: H÷"1ñ 1üò:÷<ð <ú÷	Jð Jú÷Cz	ð z	ús¦   Á
+[Ã/[ÇA\È2A\Ê
[Ê[ÊD\Î8[ÏA"\Ð(A["Ñ<B3\Ô/[/Ô<:\Õ6A\ ×A[4ØA\ Ù!A
\Û[Û
\Û"[,	Û'\Û4[=Û9\ Ü \		Ü\Ü\c                óà   — d}|j                   r|dz  }|dz  }|dz  }n|j                  s|dz  }| j                  r|j                   s
|dz  }|dz  }|j                  j	                  |d¬«       y )	Naß  
import functools
import math
import torch
import jax
import jax.numpy as jnp
from jax.experimental import pallas as pl
from torch.utils._ordered_set import OrderedSet
from torch._inductor.runtime.runtime_utils import (
    pallas_compute_tiling, pallas_make_block_spec, pallas_permute,
    pallas_gpu_align_output_specs, pallas_gpu_pad_inputs,
    pallas_gpu_unpad_results,
    pallas_ensure_nonzero_rank,
    pallas_make_block_spec_non_tiled,
    torch_dtype_to_jax_runtime,
)
z
import jax.exportz1
from jax.experimental.pallas import tpu as pltpuz8
from torch_tpu._internal.pallas import tpu_torch_pallasz8
from jax.experimental.pallas import mosaic_gpu as plgpuzA
import os as _os; _os.environ.setdefault('JAX_PLATFORMS', 'cpu')T©Ústrip)rÖ  r×  r  rÔ  Úsplice)r$   rç  Úimportss      r)   rŽ  zPallasKernel._codegen_importsò  s„   € ðˆð" �:Š:ØÐ,Ñ,ˆGØÐKÑKˆGØÐRÑR‰GØ×%Ò%ØÐRÑRˆGØ×Ò¨¯
ª
ØØTñˆGð ÐKÑKˆGØ�‰�‰˜ tˆÕ,r+   c                ó
  — d}d}t        d„ | j                  j                  «       D «       «      }| j                  j                  «       D ]8  \  }}||k(  r|j                  s|c S ||z   c S |j                  r|dz  }Œ4|dz  }Œ: y)a  Map an iteration variable to its output tensor axis index.

        Non-reduction variables map to axes 0, 1, 2, ... in order.
        Reduction variables map to axes after all pointwise axes.
        Returns None if the mapping cannot be determined.
        r   c              3  ó:   K  — | ]  }|j                   rŒd –— Œ y­wrD  r6  )r@  rª  s     r)   rA  z2PallasKernel._get_iter_var_axis.<locals>.<genexpr>  s   è ø€ ÒS˜ÀAÇNÃN”1ÑSùr7  r   N)rD  r/  rå  rN  r   )r$   rÎ  r™  Úr_idxÚn_pwrš  rÏ  s          r)   Ú_get_iter_var_axiszPallasKernel._get_iter_var_axis  s�   € ð ˆØˆÜÑS˜d×3Ñ3×:Ñ:Ó<ÔSÓSˆØ×/Ñ/×5Ñ5Ó7ò 	‰JˆC�Ø�gŠ~Ø%*×%7Ò%7�vÒI¸TÀE¹\ÒIØ×!Ò!Ø˜‘
‘à˜!‘‘ð	ð r+   c           
     óÀ  ‡— t        | j                  j                  «       D �cg c]A  }t        |j                  t
        t        j                  f«      rt        |j                  «      ‘ŒC c}«      Šˆfd„}| j                  j                  «       }|j                  | j                  j                  «       d\  }}|D ]  } ||«      }|d   sŒ|\  }} ||fS  ||fS c c}w )Nc                óþ   •— t         j                  j                  | «      }|�t        |j	                  «       «      dk  ryt        d„ |j	                  «       D «       «      }t        j                  |«      }|‰v r||fS dS )Nr   ©NNc              3  óv   K  — | ]1  }t        |t        t        j                  f«      rt        |«      n|–— Œ3 y ­wrS   )r§   rÀ  r  r†  r?  s     r)   rA  zePallasKernel._get_reshape_target_shape_and_numel.<locals>._get_nd_shape_if_matches.<locals>.<genexpr>5  s3   è ø€ ò àô % Q¬¬e¯m©mÐ(<Ô=”�A”À1ÓDñùs   ‚79)r   r·   ræ  r  r£  r¾  r©   r¡  )r{  rl  rà  r  Úiter_lengthss       €r)   Ú_get_nd_shape_if_matcheszRPallasKernel._get_reshape_target_shape_and_numel.<locals>._get_nd_shape_if_matches1  st   ø€ Ü—'‘'×(Ñ(¨Ó2ˆCØˆ{œc #§,¡,£.Ó1°QÒ6Ø!Üñ àŸ™›ôó ˆEô —I‘I˜eÓ$ˆEØ%*¨lÑ%:�E˜5�>ÐLÀÐLr+   rå  r   )r   r/  rå  r§   rO  rÀ  r  r†  rø  Úcopyr�  r#   r!  )	r$   rª  rè  Úcandidate_buf_namesÚreshape_target_shapeÚreshape_target_numelr{  r1   rç  s	           @r)   Ú#_get_reshape_target_shape_and_numelz0PallasKernel._get_reshape_target_shape_and_numel$  sä   ø€ ô
 "ð ×.Ñ.×5Ñ5Ó7öàÜ˜aŸh™h¬¬e¯m©mÐ(<Ô=ô �A—H‘H•òó
ˆô		Mð #×7Ñ7×<Ñ<Ó>ÐØ×"Ñ" 4§9¡9×#:Ñ#:Ô;à5?Ñ2ÐÐ2Ø+ò 	ˆHÙ-¨hÓ7ˆFØ�a‹yØ=CÑ:Ð$Ð&:Øà#Ð%9Ð9Ð9ð	ð $Ð%9Ð9Ð9ùò9s   £ACc                ó  — ||   }|j                   j                  }| j                  |«      }| j                  |«      }t	        |«      }| j                  |||«      }dg|z  }	||	|<   dj                  |	«      }
d|› d�}|› d|
› d�}|S )Nrò  r    rñ  r!   rs  )rÏ  rO  rÌ   rÍ   r  Ú_broadcast_axis_idxrV  )r$   Úbroadcast_varsÚbroadcast_idxÚbvrO  Úrenamed_lengthÚ
length_strÚnum_broadcast_dimsÚaxis_idxr  r  ÚarangeÚreshapeds                r)   Ú$_make_broadcasted_iteration_var_exprz1PallasKernel._make_broadcasted_iteration_var_exprH  s¨   € ð ˜MÑ*ˆØ—‘—‘ˆØ×-Ñ-¨fÓ5ˆØ—Z‘Z Ó/ˆ
ä  Ó0ÐØ×+Ñ+Ø˜MÐ+=ó
ˆð �eÐ0Ñ0ˆØ *ˆ�HÑØ—I‘I˜kÓ*ˆ	Ø˜z˜l¨!Ð,ˆØ�X˜Y y k°Ð3ˆØˆr+   c                óX  ‡— | j                   r| j                  s| j                  sy |j                  d«       | j	                  «       \  }}t        | j                   j                  «       «      }g }d }t        |«      D ]J  \  Š\  }}	| j                  |	j                  «      }
|
�|
|k(  r‰}Œ.|j                  t        ‰||	|
«      «       ŒL t        |«      }t        |«      D �]G  \  Š\  }}	|| j                  vrŒt        |«      }|	j                  }| j                  |«      }| j                  |«      }| j                  |«      }
|
€m|rS|dkD  rN‰|k7  rIt!        ˆfd„t        |«      D «       d «      }|�)| j#                  ||«      }|j                  |› d|› �«       Œº|j                  |› d|› d�«       ŒÒ|rMt        |«      dkD  r?|
|k(  r:dj%                  d„ |D «       «      }d	|› d�}|j                  |› d|› d
|› d�«       �Œ!|dkD  rL‰|k7  rGt!        ˆfd„t        |«      D «       «      }| j#                  ||«      }|j                  |› d|› �«       �Œrdj%                  d„ | j&                  j(                  D «       «      }||v }|xr: | j                   xr+ | j*                   xr | j,                   xr |	j.                   }|r| j1                  |«      nd }|�:|j                  |› d|› d|› d|› d�«       | j2                  j5                  |«       �Œ1|j                  |› d|› d�«       �ŒJ y )Nz*# Define iteration variables as JAX arraysr   c              3  óH   •K  — | ]  \  }}|j                   ‰k(  sŒ|–— Œ y ­wrS   ©rÌ  ©r@  rÀ  rŠ  rÌ  s      €r)   rA  z7PallasKernel._codegen_iteration_vars.<locals>.<genexpr>ˆ  s   øè ø€ ÒQ™t˜q !ÀAÇEÁEÈSÃLœÑQùó   ƒ"›"rr  z = jnp.arange(r!   r    c              3  ó2   K  — | ]  }t        |«      –— Œ y ­wrS   rO  r?  s     r)   rA  z7PallasKernel._codegen_iteration_vars.<locals>.<genexpr>™  s   è ø€ Ò%K°¤c¨!§fÑ%KùrB  rñ  rs  c              3  óH   •K  — | ]  \  }}|j                   ‰k(  sŒ|–— Œ y ­wrS   rü  rý  s      €r)   rA  z7PallasKernel._codegen_iteration_vars.<locals>.<genexpr>�  s$   øè ø€ ò %Ù˜!˜Q¸q¿u¹uÈ»|”Añ%ùrþ  rM  c              3  ó2   K  — | ]  }t        |«      –— Œ y ­wrS   rO  rP  s     r)   rA  z7PallasKernel._codegen_iteration_vars.<locals>.<genexpr>¬  s   è ø€ Ò(S°t¬¨T¯Ñ(SùrB  z = jnp.arange(_pallas_tile[z$]) + pl.program_id(_pallas_ax2g.get(z, 0)) * _pallas_tile[rk  )r/  rï  rÈ   rx  rí  rI  rN  rw  rG  rO  rý  rË  r  r:   rÌ   rÍ   r  rù  rV  rÐ   rZ  ró  rÿ  r   râ  rþ  r  )r$   r¡  rç  rë  rì  Ú	var_itemsrð  Útotal_var_idxrÎ  rÏ  rÑ  rõ  r_  rO  ró  rô  rñ  r%   r  r÷  r]  Úvar_in_computeÚcan_tile_relativerö  rÌ  s                           @r)   r�  z$PallasKernel._codegen_iteration_vars[  s–  ø€ ð
 ×%Ò%¨d¯kªk¸d×>QÒ>QØà×ÑÐJÔKð ×4Ñ4Ó6ñ 	3ÐÐ2ô ˜×.Ñ.×4Ñ4Ó6Ó7ˆ	àˆØˆÜ%.¨yÓ%9ò 	Ñ!ˆCÑ!�'˜5ØŸ™¨¯©Ó5ˆJØÐ%¨*Ð8LÒ*LØ #‘à×%Ñ%Ü'¨¨W°e¸ZÓHõð	ô ! Ó0Ðä%.¨yÓ%9ó H	TÑ!ˆCÑ!�'˜5Ø˜d×1Ñ1Ñ1ØÜ˜7“|ˆHØ—\‘\ˆFØ!×1Ñ1°&Ó9ˆNØŸ™ NÓ3ˆJØŸ™¨Ó/ˆJàÐ!á(Ø*¨QÒ.Ø˜}Ò,ä$(ÛQ¤y°Ó'@ÔQØó%�Mð %Ð0Ø#×HÑHØ*¨Mó ˜ð $×-Ñ-°°
¸#¸d¸VÐ.DÔEØ Ø×%Ñ%¨¨
°.ÀÀÈAÐ&NÔOØñ %ÜÐ,Ó-°Ò1ØÐ"6Ò6à ŸI™IÑ%KÐ6JÔ%KÓK�	Ø& z l°!Ð4�Ø×%Ñ%¨¨
°#°f°X¸YÀyÀkÐQRÐ&SÖTØ# aÒ'¨C°=Ò,@Ü $ó %Ü"+¨NÓ";ô%ó !�ð ×@Ñ@Ø" Mó�ð ×%Ñ%¨¨
°#°d°VÐ&<Ö=ð  $Ÿy™yÑ(S¸t¿|¹|×?RÑ?RÔ(SÓS�Ø!)¨\Ð!9�à"ò /Ø ŸK™K˜ò/à ×2Ñ2Ð2ò/ð !×5Ñ5Ð5ò/ð "×.Ñ.Ð.ð "ñ 9J�D×+Ñ+¨GÔ4Ètð ð Ð'Ø×)Ñ)Ø#˜*Ð$?À¸zð J=Ø=E¸Jð G+Ø+3¨*°Að7ôð
 ×0Ñ0×4Ñ4°WÖ=à×)Ñ)¨X¨J°nÀZÀLÐPQÐ*RÖSñQH	Tr+   c                ón   — t        d„ | D «       «      }t        d„ | D «       «      }|xr |}|r|S |dz
  |z
  S )Nc              3  ód   K  — | ](  }t        |j                  «      j                  d «      –— Œ* y­w©ÚrN©r:   rÎ  r7  ©r@  rò  s     r)   rA  z3PallasKernel._broadcast_axis_idx.<locals>.<genexpr>Ë  s(   è ø€ ò !
Ø02ŒC�—
‘
‹O×&Ñ& s×+ñ!
ùs   ‚.0c              3  óf   K  — | ])  }t        |j                  «      j                  d «       –— Œ+ y­wr  r
  r  s     r)   rA  z3PallasKernel._broadcast_axis_idx.<locals>.<genexpr>Î  s+   è ø€ ò !
Ø46”�B—J‘J“×*Ñ*¨3Ó/Ô/ñ!
ùs   ‚/1r   )rH  )rð  rñ  rõ  Úhas_reduction_varsÚhas_pointwise_varsÚis_mixeds         r)   rï  z PallasKernel._broadcast_axis_idxÂ  s\   € ô !ñ !
Ø6Dô!
ó 
Ðô !ñ !
Ø:Hô!
ó 
Ðð &Ò<Ð*<ˆÙØ Ð Ø! AÑ%¨Ñ5Ð5r+   c                óÌ  — |j                   }|j                  }|j                  }|j                  d«       |j                  d«       |j                  d«       |j                  d«       |j                  d«       |D ]  }|j                  d|› d�«       Œ |j                  d«       |j                  d	«       |j                  d
«       |D �cg c]  }|› d�‘Œ	 }}|D �cg c]  }|› d�‘Œ	 }	}|D �cg c]  }|› d�‘Œ	 }
}|D �cg c]  }|› d�‘Œ	 }}|j                  d«       |j                  d«       ||	z   }|
|z   }t	        t        |«      «      D �cg c]  }d|› �‘Œ	 }}dj                  ||z   «      }|j                  ddj                  |«      › d|› d�«       |j                  «       5  |j                  d«       |j                  d«       |j                  «       5  |j                  d«       |j                  d«       |j                  d«       t        t        ||
«      «      D ]#  \  }\  }}|j                  d|› d|› d|› d�«       Œ% |j                  d«       |j                  d«       t        |«      D ]  \  }}|j                  d|› d�«       Œ |j                  d«       |j                  d«       |
|z   }|j                  d«      j                  «       }|j                  |› ddj                  |«      › d�«       |j                  d«       |j                  d «       |j                  d!«       t        |	|«      D ]  \  }}|j                  d"|› d|› d#�«       Œ |j                  d$«       |j                  d«       |j                  d%«       d d d «       |j                  d«       |j                  d&«       |j                  d'«       d d d «       |j                  d«       |j                  d(«       |j                  d)«       t        |
«      D ]"  \  }}||   }|j                  d*|› d+|› d,�«       Œ$ t        |«      D ]  \  }}|j                  d*|› d-|› d.�«       Œ |D ]  }|j                  d*|› d/�«       Œ |j                  d«       |j                  d0«       |j                  d1«       |j                  d«       |j                  d2«       |j                  d3«       |j                  «       5  |j                  d4«       |j                  d5«       |j                  d6«       d d d «       |j                  d7«       |D ]  }|j                  d8|› d9�«       Œ |j                  d«       |j                  d«       |j                  d:«       |j                  d;«       y c c}w c c}w c c}w c c}w c c}w # 1 sw Y   �Œ,xY w# 1 sw Y   �ŒþxY w# 1 sw Y   Œ®xY w)<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Ú_smemrR  z4# Wrapper kernel using lax.fori_loop with direct TMAÚ	_barrier_r    zdef _tma_kernel(z, *, rt  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 tilesrÃ  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,), r  z*'] = plgpu.SMEM((_tile_size,), out_dtypes[r„  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,rÆ  r¾  z.flatten(),z$# Reshape results to original shapesz:return pallas_gpu_unpad_results(_result, _orig_out_shapes))rÔ  rã  rÜ  rx  rK  r  rV  r�  rw  r\  ÚrstriprÛ  )r$   rç  rµ  rÔ  rã  rÜ  rz  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_inrE  Úkernel_call_argsrJ   Úgmem_outÚsmem_outÚ
smem_paramÚ
orig_paramÚbarrier_params                              r)   r˜  z%PallasKernel._codegen_jit_wrapper_tmaÖ  sÕ  € Ø�x‰xˆØ!×5Ñ5ÐØ×)Ñ)ˆð 	�‰ÐOÔPØ�‰Ð,Ô-Ø�‰Ð;Ô<Ø�‰Ð6Ô7à�‰Ð'Ô(Ø(ò 	JˆEØ�N‰NÐ:¸5¸'ÀÐHÕIð	Jà�‰Ð1Ô2Ø�‰ÐKÔLà�‰ÐQÔRà2EÖF¨Q ˜s %š[ÐFÐÐFØ3@ÖA¨a   5škÐAÐÐAØ2EÖF¨Q ˜s %š[ÐFÐÐFØ3@ÖA¨a   5škÐAÐÐAà�‰�rÔØ�‰ÐMÔNà 1Ð4FÑ FÐØ+Ð.@Ñ@ˆÜ38¼Ð=PÓ9QÓ3RÖS¨a˜I a Sš/ÐSˆÐSØŸ™ ?°^Ñ#CÓDˆà�‰Ø˜tŸy™yÐ)>Ó?Ð@ÀÀnÐEUÐUWÐXô	
ð �[‰[‹]ñ )	MØ�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ò H‘D�A�qØ—N‘NÐ%AÀ!ÀÀAÐ#FÕGðHð —‘˜rÔ"Ø—‘Ð8Ô9Ø#4Ð7IÑ#IÐ Ø&×-Ñ-¨cÓ2×8Ñ8Ó:�	Ø—‘ ) ¨A¨d¯i©iÐ8HÓ.IÐ-JÈ!ÐLÔMà—‘˜rÔ"Ø—‘ØMôð —‘Ð4Ô5Ü*-Ð.@ÐBTÓ*Uò Ñ&�H˜hØ—N‘NØ2°8°*¸B¸x¸jÐHlÐmõðð —‘Ð;Ô<Ø—‘˜rÔ"Ø—‘˜}Ô-÷E".ðH �N‰N˜2ÔØ�N‰NÐ5Ô6Ø�N‰NÐKÔL÷S)	MðX 	�‰�rÔØ�‰ØOô	
ð 	�‰Ð-Ô.Ü&Ð'8Ó9ò 	‰MˆAˆzØ,¨QÑ/ˆ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ÔØ(ò 	6ˆEØ�N‰N˜T % ¨Ð4Õ5ð	6à�‰�sÔà�‰�rÔØ�‰Ð=Ô>Ø�‰ÐSÕTùòG GùÚAùÚFùÚAùò T÷".ñ ".ú÷)	Mñ )	Mú÷P	>ð 	>úsO   ÃV'ÃV,Ã1V1ÄV6ÅV;Æ,3WÇF7W Î;WÓ94W× W
	×W×W×W#c                ó¸  — |j                   }|j                  }ddj                  |«      › d�}|j                  d«       |j                  d«       |D ]  }|j                  d|› d�«       Œ |j                  d«       |j                  d	«       |j                  d
«       |j                  d«       |j                  d«       |j                  d«       |j                  d«       |j                  d|› d�«       |j                  d«       |j                  d«       |j                  d|z   «       |j                  d«       |j                  d«       |j                  d«       |j                  d«       |j                  d«       |j                  d«       |j                  d«       |j                  d«       |j                  d«       |j                  d|› d�«       |j                  d«       |j                  d«       |j                  d «       |j                  d!«       |j                  d"«       |j                  d#|z   «       |j                  d$«       |j                  d%«       |j                  d&«       |j                  d'«       |j                  d(«       |j                  d)«       |j                  d*«       |j                  d+|› �«       |j                  d,«       |j                  d-«       |j                  d.«       |j                  d"«       |j                  d#|z   «       |j                  d$«       |j                  d%«       |j                  d/«       y )0Nrj  r    rk  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 = OrderedSet(_all_sizes)zH_can_pad = len(_unique_sizes) == 1 and all(s > 1 for s in _unique_sizes)rR  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)r¿  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))rÔ  rã  rV  rx  )r$   rç  rµ  rÔ  rã  Ú
input_listrz  s          r)   r™  z,PallasKernel._codegen_jit_wrapper_legacy_gpuN  sä  € ð �x‰xˆØ!×5Ñ5ÐØ˜Ÿ™Ð#6Ó7Ð8¸Ð:ˆ
ð
 	�‰ÐPÔQØ�‰Ð(Ô)Ø(ò 	?ˆEØ�N‰NÐ/°¨w°fÐ=Õ>ð	?à�‰Ð1Ô2Ø�‰Ð@ÔAØ�‰Ð?Ô@Ø�‰Ø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                ón   — | j                   j                  j                  «       D ]  \  }}||k(  sŒ|c S  y)z:Map a kernel parameter name back to its graph buffer name.N)r#   r!  rN  )r$   rz  Ú
graph_nameÚ
inner_names       r)   rv  zPallasKernel._param_to_buf_nameœ  s<   € à&*§i¡i×&=Ñ&=×&CÑ&CÓ&Eò 	"Ñ"ˆJ˜
Ø˜UÓ"Ø!Ò!ð	"ð r+   c                ó  — |j                   }| j                  }| j                  rdnd}|j                  rdnd}g }|j                  D ]A  }| j                  |«      }|j                  |r| j                  j                  |«      nd«       ŒC | j                  }	|	rd|	› �nd}
t        d„ |D «       «      rdt        |«      › �nd}|j                  d|› d	|› d
|› |› |
› d�	«       |j                  d«       |j                  d«       |j                  d«       |j                  d«       |j                  d«       |j                  d«       | j                  ||j                  «       |j                  D ]O  }| j                  |«      }|r| j                  j                  |«      nd}|€Œ6|j                  |› d|› d|› d�«       ŒQ dj                  |j                  «      }t        d„ |D «       «      redj                  d„ |D «       «      }|j                  d|› d�«       |j                  d«       |j                  d|› d�«       |j                  d«       y|j                  d«       |j                  d|› d�«       |j                  d«       y)uQ  Generate tiled BlockSpec and grid variables for CPU/TPU.

        Tiles the last 1â€“2 dimensions of each tensor, respecting TPU
        alignment constraints (last dim multiple of 128, second-to-last
        multiple of 8).  Lower-ndim inputs are right-aligned with the
        reference output shape per numpy broadcast rules.
        rÙ   rÚ   Nz, max_grid_product=rR  c              3  ó$   K  — | ]  }|d u–— Œ
 y ­wrS   r?   ©r@  r'   s     r)   rA  z4PallasKernel._codegen_tiled_specs.<locals>.<genexpr>¼  s   è ø€ Ò4 Q�1˜D”=Ñ4ùrš  z, permutations=zM_tile, _grid, _ax2g = pallas_compute_tiling(_pallas_out_shapes[0], transpose=z, skip_last_n=z4, exact_only=len(_pallas_out_shapes[0]) < 2, is_tpu=r!   z_ng = len(_grid)z_ref = _pallas_out_shapes[0]r…  zF    pallas_make_block_spec(s, _ref, _tile, _ax2g, _ng, is_output=True)z    for s in _pallas_out_shapesrr  rs  r    c              3  ó$   K  — | ]  }|d u–— Œ
 y ­wrS   r?   r.  s     r)   rA  z4PallasKernel._codegen_tiled_specs.<locals>.<genexpr>Ü  s   è ø€ Ò0 ˆq˜Œ}Ñ0ùrš  c              3  ó2   K  — | ]  }t        |«      –— Œ y ­wrS   )r¬   r.  s     r)   rA  z4PallasKernel._codegen_tiled_specs.<locals>.<genexpr>Ý  s   è ø€ Ò!=¨a¤$ q§'Ñ!=ùrB  z_perm_flags = [rk  r†  z]    pallas_make_block_spec(i.shape, _ref, _tile, _ax2g, _ng, permutation=p) for i, p in zip([z], _perm_flags)zG    pallas_make_block_spec(i.shape, _ref, _tile, _ax2g, _ng) for i in [)rÔ  r[  rZ  rÖ  rã  rv  rý  rô  r»  r÷  rH  r¬   rx  r€  rõ  rV  )r$   rç  rÔ  Úskip_nÚhas_transposeÚis_tpu_literalÚ	all_permsr'   r{  ÚmgpÚmgp_argÚ	perms_argrz  r²  r(  Ú	perm_lists                   r)   r—  z!PallasKernel._codegen_tiled_specs£  s˜  € ð �x‰xˆØ×&Ñ&ˆØ"&×"9Ò"9™¸wˆØ#&§:¢:™°7ˆð 35ˆ	Ø×(Ñ(ò 	ˆAØ×.Ñ.¨qÓ1ˆHØ×ÑÙ=E�×+Ñ+×/Ñ/°Ô9È4õð	ð ×(Ñ(ˆÙ14Ð'¨ uÑ-¸"ˆô Ñ4¨)Ô4Ô4ð œd 9›oÐ.Ñ/àð 	ð
 	�‰ðà&˜ð (Ø!˜(ð #Ø$Ð%Øˆk˜'˜ !ð%ô	
ð 	�‰Ð)Ô*Ø�‰Ð5Ô6à�‰Ð2Ô3Ø�‰ØTô	
ð 	�‰Ð8Ô9Ø�‰�sÔà×&Ñ& t¨S×-DÑ-DÔEð ×,Ñ,ò 	GˆEØ×.Ñ.¨uÓ5ˆHÙDL�T×2Ñ2×6Ñ6°xÔ@ÐRVˆFØÑ!Ø—‘ % ¨¨E¨7°)¸F¸8À1ÐEÕFð		Gð —Y‘Y˜s×6Ñ6Ó7ˆ
ÜÑ0 iÔ0Ô0ØŸ	™	Ñ!=°9Ô!=Ó=ˆIØ�N‰N˜_¨Y¨K°qÐ9Ô:Ø�N‰NÐ5Ô6Ø�N‰Nð%Ø%/ L°ðAôð �N‰N˜3Õà�N‰NÐ5Ô6Ø�N‰NØYÐZdÐYeÐefÐgôð �N‰N˜3Õr+   c                ó  — |j                   }| j                  rdnd}|j                  d«       |j                  d|z   «       |j                  d«       |j                  d«       |j                  d«       |j                  d|j                  › d	�«       |j                  d
|› d	�«       |j                  |rd|› d�nd«       |j                  d«       |j                  rA|j                  D �cg c]  }d|› d�‘Œ
 }}|j                  ddj                  |«      › d	�«       |j                  d«       |j                  d«       |j                  d«       |j                  d«       |j                  d«       |j                  d«       y c c}w )NÚ_gridz(1,)rÂ  r¾  z     out_shape=out_shapes_pallas,z    out_specs=out_specs_pallas,z    in_specs=in_specs_pallas,z    interpret=rÃ  z	    grid=z    input_output_aliases={ rÄ  z    input_output_aliases={},rÆ  zpallas_ensure_nonzero_rank(r!   r    zif isinstance(_result, tuple):zF    _result = tuple(r.reshape(s) for r, s in zip(_result, out_shapes))r¿  z,    _result = _result.reshape(out_shapes[0])zreturn _result)rÔ  r‘  rx  rØ  rã  rV  )	r$   rç  rµ  ré  r­  rÔ  Ú	grid_exprr'   Ú kernel_input_params_nonzero_ranks	            r)   rš  z)PallasKernel._codegen_jit_wrapper_cpu_tpuì  sm  € ð �x‰xˆØ#×0Ò0‘G°fˆ	Ø�‰Ð2Ô3Ø�‰�v 
Ñ*Ô+Ø�‰Ð9Ô:Ø�‰Ð8Ô9Ø�‰Ð6Ô7Ø�‰˜¨×(=Ñ(=Ð'>¸aÐ@ÔAØ�‰˜ 9 +¨QÐ/Ô0Ø�‰áð +Ð+<Ð*=¸TÑBà/ô	
ð
 	�‰�tÔØ×"Ò"à<?×<SÑ<Sö0Ø78Ð-¨a¨S°Ò2ð0Ð,ð 0ð �N‰N˜T $§)¡)Ð,LÓ"MÐ!NÈaÐPÔQØ�‰�sÔà�‰Ð7Ô8Ø�‰ØTô	
ð 	�‰�wÔØ�‰ÐEÔFØ�‰Ð'Õ(ùò0s   Ã-Fc                óf   — |j                   r| j                  ||«       y | j                  ||«       y rS   )rÖ  Ú_codegen_main_entry_tpuÚ_codegen_main_entry_default)r$   rç  r¦  s      r)   r›  z PallasKernel._codegen_main_entry  s+   € Ø�:Š:Ø×(Ñ(¨Ð.>Õ?à×,Ñ,¨SÐ2BÕCr+   c                ó
  — |j                   }|j                  d«       |j                  › d�}|j                  }|j                  d|› ddj                  |j                  «      › d�«       |j                  «       5  |j                  d«       |j                  d«       t        |j                  «      t        |j                  «      z   }|D ]  }|j                  |› d	|› d
|› d|› �«       Œ  |j                  d«       g }|j                  D ]0  }	|j                  |	› d|	› d|	› d�«       |j                  |	› d�«       Œ2 |j                  D ]0  }
|j                  |
› d|
› d|
› d�«       |j                  |
› d�«       Œ2 |j                  ddj                  |j                  D �cg c]  }d|› d�‘Œ
 c}«      z   dz   «       g }|j                  D ]p  }|j                  j                  |«      }|�<t        j                  j                  |«      }|�|j                  t!        |«      «       Œ\|j                  d|› d�«       Œr |j                  ddj                  |«      z   dz   «       ddg}|j#                  |j$                  «       |j#                  |«       |j                  d|› ddj                  |«      › d�«       g }|j                  D ]  }|j                  d|› d�«       Œ |rdj                  |«      nd}g }|j&                  D ]  }|j                  d|› d�«       Œ |rdj                  |«      nd}|j                  d |› d!|› d"|› �«       |j                  d#|› d$�«       |j                  «       5  |j                  d%«       |j                  «       5  |j                  d&|› d'�«       d d d «       |j                  d(«       |j                  «       5  |j                  d&|› d)�«       d d d «       d d d «       t        |j                  «      t        |j                  «      z   }|j                  d*dj                  |«      › d+�«       |j                  d,«       |j                  «       5  |j                  D ]q  }|j                  j                  |«      }|�:t        j                  j                  |«      }|�|j                  d-|› d.|›d/�«       ŒZ|j                  d-|› d.|› d0�«       Œs 	 d d d «       |j                  d+«       |j(                  rodj                  d1„ |j(                  D «       «      }|j                  d2|› d3�«       dj                  d4„ |j(                  D «       «      }|j                  d5|› d+�«       n"|j                  d6«       |j                  d7«       |j                  d%«       |j                  «       5  |j                  d8|› d9�«       d d d «       |j                  d(«       |j                  «       5  |j                  d8|› d:�«       d d d «       |j                  d;«       |j                  «       5  |j                  d<«       |j                  «       5  |j                  d=«       d d d «       d d d «       d d d «       y c c}w # 1 sw Y   �ŒêxY w# 1 sw Y   �Œ¸xY w# 1 sw Y   �Œ½xY w# 1 sw Y   �ŒÐxY w# 1 sw Y   ŒïxY w# 1 sw Y   Œ¼xY w# 1 sw Y   ŒkxY w# 1 sw Y   ŒoxY w# 1 sw Y   y xY w)>NrR  Ú_mainrr  ré   r    ú, stream=None):z*jax.config.update('jax_enable_x64', False)újax.clear_caches()rr  z.to(torch.int32) if z.dtype == torch.int64 else z+# Build JAX placeholders for export tracingz$_placeholder = jax.ShapeDtypeStruct(z#.shape, torch_dtype_to_jax_runtime(z.dtype))Ú_placeholderúout_shapes = (útuple(rÇ  rx  útorch_dtype_to_jax_runtime(r  úout_dtypes = (r|  r}  zexported = jax.export.export(z, platforms=['tpu'])(r!   z'_'.join(str(s) for s in z	 + 'x' + z''zkernel_key = 'z	_out_' + z + '_in_' + z.if not tpu_torch_pallas.lookup_custom_kernel('z', kernel_key):ztry:z)tpu_torch_pallas.register_custom_kernel('zF', kernel_key, serialized_mlir_module=exported.mlir_module_serialized)zexcept TypeError:z/', kernel_key, exported.mlir_module_serialized)zinput_tensors = [rk  zoutput_shape_tensors = [ztorch.empty(z.shape, dtype=z, device='tpu'),z.dtype, device='tpu'),c              3  ó0   K  — | ]  \  }}|› d |› �–— Œ y­wr  r?   r�  s      r)   rA  z7PallasKernel._codegen_main_entry_tpu.<locals>.<genexpr>   s   è ø€ Ò)T¹&¸1¸a¨Q¨C¨r°!°¬+Ñ)Tùrƒ  z_input_output_aliases = { z }c              3  ó8   K  — | ]  \  }}t        |«      –— Œ y ­wrS   rO  )r@  rÀ  Ú_os      r)   rA  z7PallasKernel._codegen_main_entry_tpu.<locals>.<genexpr>¢  s   è ø€ Ò.U¹'¸1¸b¬s°1¯vÑ.Uùs   ‚z_donate_argnums = [z_input_output_aliases = {}z_donate_argnums = []z0_results = tpu_torch_pallas.call_custom_kernel('zh', kernel_key, inputs=input_tensors, output_shapes=output_shape_tensors, donate_argnums=_donate_argnums)zt', kernel_key, inputs=input_tensors, output_shapes=output_shape_tensors, input_output_aliases=_input_output_aliases)zif _results is not None:z7for _in_idx, _out_idx in _input_output_aliases.items():z0input_tensors[_in_idx].copy_(_results[_out_idx]))rÔ  rx  rÕ  rV  rä  r�  rI  rá  râ  rý  rÜ  rÞ  r»  r   r·   r  r   r�  rÝ  rã  ré  )r$   rç  r¦  rÔ  Ú	main_nameÚkernel_name_strÚall_input_paramsÚ
param_nameÚall_jax_input_namesr¹  Úptrr£  Údtype_exprsr{  r»   Úwrapper_placeholder_argsÚshape_key_partsr'   Úoutput_key_exprÚinput_key_partsÚinput_key_exprÚinput_tensor_namesÚalias_map_strÚdonate_argnums_strs                           r)   r>  z$PallasKernel._codegen_main_entry_tpu  sb  € ð �x‰xˆØ�‰�rÔØ—‘Ð' uÐ-ˆ	ØŸ/™/ˆØ�‰Ø�9�+˜Q˜tŸy™y¨×)?Ñ)?Ó@ÐAÀÐQô	
ð �[‰[‹]ñ d	Wð
 �N‰NÐGÔHØ�N‰NÐ/Ô0ô  $ C×$4Ñ$4Ó5¼¸S×=MÑ=MÓ8NÑNÐØ.ò �
Ø—‘Ø!�l # j \ð 2Ø$˜Ð%@ÀÀðNõðð �N‰NÐHÔIØ"$ÐØ!×.Ñ.ò H�
Ø—‘Ø!�lÐ"FØ!�lÐ"EÀjÀ\ÐQYð[ôð $×*Ñ*¨j¨\¸Ð+FÕGðHð ×'Ñ'ò A�Ø—‘Ø�eÐ?Ø�eÐ>¸s¸eÀ8ðMôð $×*Ñ*¨c¨U°,Ð+?Õ@ðAð �N‰NØ Ø—)‘)À×@QÑ@QÖR¸˜v d V¨7Ò3ÒRÓSñTàñôð
 &(ˆKØ×)Ñ)ò P�Ø×3Ñ3×7Ñ7¸Ó=�ØÐ'ÜŸG™G×-Ñ-¨hÓ7�EØÐ(Ø#×*Ñ*Ô+=¸eÓ+DÔEØ Ø×"Ñ"Ð%@ÀÀÀgÐ#NÕOðPð �N‰NÐ+¨d¯i©i¸Ó.DÑDÀtÑKÔLð )5°lÐ'CÐ$Ø$×+Ñ+¨C×,?Ñ,?Ô@Ø$×+Ñ+Ð,?Ô@Ø�N‰NØ/Ø#Ð$ð %Ø—I‘IÐ6Ó7Ð8¸ð;ôð !ˆOØ×&Ñ&ò O�Ø×&Ñ&Ð)BÀ1À#ÀWÐ'MÕNðOñ 6E�× Ñ  Ô1È$ð ð !ˆOØ×,Ñ,ò O�Ø×&Ñ&Ð)BÀ1À#ÀWÐ'MÕNðOñ 6E�× Ñ  Ô1È$ð ð �N‰NØ  Ð 1°Ø"Ð#Ø˜~Ð.ð0ôð �N‰NØ@ÀÐ@QÐQ`Ðaôð —‘“ñ Ø—‘˜vÔ&Ø—[‘[“]ñ Ø—N‘NðØ+Ð,ð -RðSô÷ð —‘Ð2Ô3Ø—[‘[“]ñ Ø—N‘NðØ+Ð,ð -;ð<ô÷÷ô" "& 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Ôð �ŠØ $§	¡	Ñ)TÀCÇOÁOÔ)TÓ T�Ø—‘Ð!<¸]¸OÈ3ÐOÔPØ%)§Y¡YÑ.UÀSÇ_Á_Ô.UÓ%UÐ"Ø—‘Ð!4Ð5GÐ4HÈÐJÕKà—‘Ð;Ô<Ø—‘Ð5Ô6à�N‰N˜6Ô"Ø—‘“ñ Ø—‘ðØ'Ð(ð )6ð7ô÷ð �N‰NÐ.Ô/Ø—‘“ñ Ø—‘ðØ'Ð(ð )BðCô÷ð �N‰NÐ5Ô6Ø—‘“ñ WØ—‘ØMôð —[‘[“]ñ WØ—N‘NÐ#UÔV÷W÷	W÷d	Wð d	WùòF S÷fñ ú÷ñ ú÷ñ ú÷,ñ ú÷<ð ú÷ð ú÷$Wð Wú÷	Wð Wú÷d	Wð d	Wúsã   Á:D'[9Æ!ZÆ.G[9Í5"Z/ÎZÎ-)Z/ÏZ"Ï,Z/Ï4A8[9Ñ,BZ<Ó.C[9×[	×)[9Ø[Ø)[9Ù"[-Ù%[!Ù7[-Ù?[9Ú[9ÚZÚZ/Ú"Z,Ú'Z/Ú/Z9	Ú4[9Ú<[	Û[9Û	[	Û[9Û[	Û[9Û![*Û&[-Û-[6	Û2[9Û9\c                ó†  ‡— |j                   }|j                  d«       |j                  › d�}|j                  d|› ddj                  |j                  «      › d�«       |j                  «       5  |j                  d«       |j                  r|j                  d«       |j                  d	«       |j                  rK|j                  d
«       |j                  D ]+  }| j                  |||j                  |j                  ¬«       Œ- |j                  d«       |j                  D ]3  }|j                  d«      sŒ| j                  |||j                  d¬«       Œ5 |j                  d«       |j                  D ]3  }|j                  d«      sŒ| j                  |||j                  d¬«       Œ5 |j                  d«       |j                  ddj                  |j                  D �cg c]  }d|› d�‘Œ
 c}«      z   dz   «       g }|j                  D ]p  }|j                  j                  |«      }	|	�<t        j                   j#                  |	«      }
|
�|j%                  t'        |
«      «       Œ\|j%                  d|› d�«       Œr |j                  ddj                  |«      z   dz   «       i Š|j                  D ]
  }|› d�‰|<   Œ |j                  D ]
  }|› d�‰|<   Œ ddg}|j)                  |j*                  «       |j)                  ˆfd„|j,                  D «       «       |j                  d|› ddj                  |«      › d�«       |j                  d «       |j.                  rH|j                  d!«       |j.                  D ](  }|j                  |   }|j                  |› d"|› d#�«       Œ* d d d «       y c c}w # 1 sw Y   y xY w)$NrR  rA  rr  ré   r    rB  z)jax.config.update('jax_enable_x64', True)z>jax.config.update('jax_default_device', jax.devices('cpu')[0])rC  z*# Convert Torch -> JAX for donated outputs)Ú
contiguousz+# Convert Torch -> JAX for in-place tensorsrn  Fz!# Convert Torch -> JAX for inputsrq  Tz-# Prepare output metadata from PyTorch tensorrE  rF  rÇ  rx  rG  r  rH  Ú_jaxr|  r}  c              3  ó(   •K  — | ]	  }‰|   –— Œ y ­wrS   r?   )r@  r£  Úarg_name_maps     €r)   rA  z;PallasKernel._codegen_main_entry_default.<locals>.<genexpr>  s   øè ø€ ò %Ø'+�˜TÕ"ñ%ùrò  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])))rÔ  rx  rÕ  rV  rä  r�  r×  rá  Ú_emit_torch_to_jaxrÖ  râ  r7  rÜ  rÞ  r»  r   r·   r  rý  r   r�  rÝ  rã  rç  )r$   rç  r¦  rÔ  rL  r¹  rQ  r£  rR  r{  r»   Úwrapper_call_argsrÌ  r¬  r_  s                 @r)   r?  z(PallasKernel._codegen_main_entry_defaultÇ  sÎ  ø€ ð �x‰xˆØ�‰�rÔØ—‘Ð' uÐ-ˆ	Ø�‰Ø�9�+˜Q˜tŸy™y¨×)?Ñ)?Ó@ÐAÀÐQô	
ð �[‰[‹]ñ E	Ø�N‰NÐFÔGØ×#Ò#Ø—‘ØTôð �N‰NÐ/Ô0Ø×ÒØ—‘ÐKÔLØ"%×"2Ñ"2ò �Jð ×+Ñ+ØØ"ØŸ
™
Ø#&×#7Ñ#7ð	 ,õ ðð �N‰NÐHÔIØ×'Ñ'ò U�Ø—>‘> ,Õ/Ø×+Ñ+¨D°#°s·z±zÈeÐ+ÕTðUð �N‰NÐ>Ô?Ø×'Ñ'ò T�Ø—>‘> (Õ+Ø×+Ñ+¨D°#°s·z±zÈdÐ+ÕSðTð �N‰NÐJÔKØ�N‰NØ Ø—)‘)À×@QÑ@QÖR¸˜v d V¨7Ò3ÒRÓSñTàñôð
 &(ˆKØ×)Ñ)ò P�Ø×3Ñ3×7Ñ7¸Ó=�ØÐ'ÜŸG™G×-Ñ-¨hÓ7�EØÐ(Ø#×*Ñ*Ô+=¸eÓ+DÔEØ Ø×"Ñ"Ð%@ÀÀÀgÐ#NÕOðPð �N‰NÐ+¨d¯i©i¸Ó.DÑDÀtÑKÔLØ+-ˆLØ!×.Ñ.ò ?�
Ø.8¨\¸Ð+>�˜ZÒ(ð?à×'Ñ'ò 1�Ø'* e¨4 L�˜SÒ!ð1ð ".¨|Ð <ÐØ×$Ñ$ 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õð÷CE	ð E	ùòF S÷GE	ð E	ús.   Á/CN7Ä8AN7ÆAN7Ç"N2Ç/F:N7Î2N7Î7O c               óB   — |rdnd}| j                  |› d|› |› d�«       y )Nz.detach().contiguous()z	.detach()z_jax = jax.dlpack.from_dlpack(r!   )rx  )rÔ  r_  rÖ  r\  Úsuffixs        r)   r`  zPallasKernel._emit_torch_to_jax  s-   € ñ .8Ñ)¸[ˆØ�‰˜(˜Ð#AÀ(ÀÈFÈ8ÐSTÐUÕVr+   c                ó  — t         j                  j                  }| j                  j	                  «       \  }}}}|D �cg c]  }|j
                  ‘Œ }}|D �	cg c]  }	|	j                  d«      sŒ|	‘Œ }
}	t        t        t        |«      «      }t        | di «      }|
D �	cg c](  }	|j                  |	d«      r||j                  |	«         ‘Œ* }}	|› ddj                  ||z   «      › d�}|j                  |«       yc c}w c c}	w c c}	w )z7Generate the Python code that calls this Pallas kernel.rm  r�  Fz.run(r    r!   N)r   r·   Úwrapper_coder#   r‹  r£  r7  rI  r‘  r:   r…  r»  r  rV  rx  )r$   r£  r“  Úwrapperrœ  r�  rE  r˜   Úkernel_param_namesr'   rÛ  Úcall_arg_strsÚ	aliasableÚalias_call_argsÚkernel_calls                  r)   Úcall_kernelzPallasKernel.call_kernel  sþ   € ä—'‘'×&Ñ&ˆØ$(§I¡I×$<Ñ$<Ó$>Ñ!ˆ�)˜Q Ø.6Ö7¨˜aŸf›fÐ7ÐÐ7Ø&8ÖT ¸A¿L¹LÈÕ<Sš1ÐTˆÐTÜœS¤ iÓ0Ó1ˆÜ˜DÐ"6¸Ó;ˆ	ð %ö
àØ�}‰}˜Q Ô&ð Ð,×2Ñ2°1Ó5Ó6ð
ˆð 
ð ˜˜e D§I¡I¨oÀÑ.MÓ$NÐ#OÈqÐQˆØ×Ñ˜+Õ&ùò 8ùÚTùò
s   ¾C9ÁC>Á.C>Â-D)
r%   r8   r
  r8   r  rÛ   r  rÛ   r9   ÚNone)r  r8   r9   r:   )r9   r   )r  r8   r9   r   )r  r8   r9   rÛ   )r  r8   r9   zlist[sympy.Symbol])r  r8   r£  r:   r9   z!list[tuple[int, int, int]] | None)rl  r:   rm  zlist[tuple[int, int, int]]r9   r:   )rÔ  r   ry  rÙ  r9   rm  )rO  rÀ  r9   úlist[int] | None)r�  rI  r9   zdict | None)r£  r:   r  r8   r9   ztuple[int, ...] | None)r£  r:   r  r8   r9   z.tuple[tuple[int, ...], tuple[int, ...]] | None)rÐ  ræ  r½  ræ  r9   rn  )r  r8   r9   rÇ  )rÜ   r   r9   rÐ  )rà  ræ  r9   ræ  )r•  ræ  r|  ræ  r9   rn  )rç  rÓ  r9   ztuple[bool, bool])r9   útuple[int, ...])r  r   r9   rÐ  )r9   rÐ  )r9   rÛ   )r£  r:   r9   z'tuple[Any, Any, Any, list, bool] | None)r  r8   r9   ztuple[int, OrderedSet])r  r8   r$  r   r9   r   )
rP  rI  r  rI  r¦  rÛ   r6  r   r9   rÛ   )r£  r:   r  r8   rE  rÇ  r9   rÇ  )rV  r:   r¯  ro  r9   r:   )r_  r:   r9   rT   )rl  r:   r£  r:   r  r8   r9   z_IndirectAccessInfo | None)r9   rm  )
rl  r:   r£  r:   r  r8   rE  rÇ  r9   r:   )r£  r:   rV  r:   r9   r:   )r£  r:   r  r8   rV  r:   r9   r:   )r  r8   rE  rÇ  r9   rÇ  )r{  r:   r  r8   r¤  r   r9   rÛ   )r£  r:   r9   rÛ   )r¾  r:   r¿  r   rÀ  rÛ   r9   rÙ  rS   )r¾  r:   r£  r:   r  r8   r¿  r   rE  rÇ  rÉ  r   r9   rÙ  )r¾  r:   r¿  r   r×  zdict[str, Any]r£  r:   rÉ  r   r9   r:   )r£  r:   r  r8   r9   r   )
r£  r:   r  r8   r¿  r   rÉ  r   r9   rm  rú  )r  r8   rÔ   rÍ  rí  úint | floatr9   rp  )rR  )r  r8   r!  r:   r9   údict[str, Any] | None)r!  r:   rÂ  r:   rû  rÀ  r9   rq  )r  r8   rÂ  r:   rû  rÀ  r9   rq  )
r»   r³  r¼   r³  r<  rC   r¿  ú%CSEVariable | tuple[CSEVariable, ...]r9   rr  )rJ  r:   r9   rÛ   )r£  rT   r9   r:   )rç  rÓ  rà  rß  r9   rè  )rÔ  r   r¡  r   rç  rÓ  r9   rm  )rç  rÓ  rÕ  r:   r¡  r   r9   rm  )rç  rÓ  r9   rm  )rÎ  rÍ  r9   rÐ  )r9   z)tuple[tuple[int, ...] | None, int | None])rð  úlist[_BroadcastedIterVar]rñ  rÀ  r9   r:   )r¡  r   rç  rÓ  r9   rm  )rð  rs  rñ  rÀ  rõ  rÀ  r9   rÀ  )rç  rÓ  rµ  r:   r9   rm  )rz  r:   r9   rT   )
rç  rÓ  rµ  r:   ré  rè  r­  r:   r9   rm  )rç  rÓ  r¦  r:   r9   rm  )
rÔ  r   r_  r:   rÖ  rÛ   r\  rÛ   r9   rm  )r£  r:   r“  zIRNode | Noner9   rm  )Zr;   r<   r=   r>   rX   Ú	overridesÚpallas_pexprrÍ   rÅ  rN   r  r  r  r-  r+  rÊ   r4  r9  r;  rg  r´  rp  r€  r‹  r—  rµ  rÌ  rº  rÝ  rG  rJ  r¤  rê  r  r  rû  r*  rF  rM  r8  r?  rM  rO  rT  rW  r^  rr  r~  r‰  r�  rœ  r§  rŸ  rº  rÂ  rÌ  rä  Útyping_extensionsÚoverrideré  rØ  r  r¥  r  r  r   rH  rK  rk  r·  r’  r•  r“  rŽ  râ  rí  rù  r�  rï  r˜  r™  rv  r—  rš  r›  r>  r?  r`  rl  Ú__classcell__©r  s   @r)   rë  rë  ™  së  ø… ñð &€IØ)5€EÐ&Ó5ô.+ð`8Øð8Ø&0ð8Ø9=ð8ØFJð8à	ó8ó!>óFbóHó<8ó:ó5óKó7ðIØðIØ'*ðIà	*óIðV ò	,ó ð	,ð(SØ"ð(SØ,5ð(Sà	ó(SðT òó ðó  ðDhØðhØ *ðhà	óhðTrØðrØ *ðrà	7órðh òó ðó,Uð0 òó ðð òó ðð òó ðó*&ó dóL	ó	óCóJ(KðT'Øð'à	ó'ð(ØðØ,6ðà	óð #àð#ð ð#ð ð	#ð
 !ð#ð 
ó#ðJ8àð8ð ð8ð "ð	8ð
 
ó8ðt4àð4ð ð4ð "ð	4ð
 
ó4ðlHIàðHIð ðHIð "ð	HIð
 
óHIðT ò6ó ð6óð$I
ØðI
Ø!ðI
Ø*4ðI
à	#óI
óV(
ðT8àð8ð ð8ð ð	8ð
 "ð8ð 
ó8ótð4YKØðYKØ *ðYKØ7:ðYKà	óYKðv>Øð>Ø+:ð>à	ó>ð@&9Øð&9Ø)3ð&9ØEOð&9à	ó&9óP2ðhØðØ*ðØ=Aðà	óð@ ð7:àð7:ð ð7:ð ð	7:ð
 ð7:ð "ð7:ð ð7:ð 
ó7:ðrN
àðN
ð ðN
ð %ð	N
ð
 ðN
ð ðN
ð 
óN
ð` ×Ñò-ó  ð-ó^@ðD ×ÑàLPð:7Øð:7Ø *ð:7Ø3>ð:7ØFIð:7à	ò:7ó  ð:7ðx àEFð
Øð
Ø ,ð
Ø7Bð
à	ò
ó ð
ð 57ðNØðNØ.1ðNà	óNð*
Øð
Ø.1ð
ØCFð
à	ó
ðB+
Øð+
Ø/2ð+
ØDGð+
à	ó+
ðZmàðmð ðmð &ð	mð
 5ðmð 
/ómð^ ò&ó ð&ó
DôLQðf ðØðØ/>ðà	òó ðð+àð+ð $ð+ð ð	+ð
 
ó+ð l8àðl8ð ðl8ð $ð	l8ð
 
ól8ó\-ó>ð&":à	2ó":ðHØ7ðØHKðà	óð&eTØ)ðeTØ0?ðeTà	óeTðN ð6Ø1ð6àð6ð  ð6ð 
ò	6ó ð6ó&vUðpL
Ø"ðL
Ø03ðL
à	óL
ó\óG ðR#)àð#)ð ð#)ð +ð	#)ð
 ð#)ð 
ó#)óJDðnWØ"ðnWØ69ðnWà	ónWð`NØ"ðNØ69ðNà	óNð` ðWØðWØ(+ðWØ59ðWØJNðWà	òWó ðW÷'ð 'r+   rë  c                  óT   ‡ — e Zd ZeZedd„«       Zˆ fd„ZeZeZ		 	 	 	 	 	 	 	 dd„Z
ˆ xZS )ÚPallasSchedulingc                ó6   — t        t        j                  g«      S rS   )r   r   ÚREDUCE_TO_SINGLE_ELEMENT)Úclsr  s     r)   Úget_backend_featuresz%PallasScheduling.get_backend_features6  s   € ô œ>×BÑBÐCÓDÐDr+   c                ó¤  •— t         ‰| �  ||«      sy|j                  «       r®|j                  «       ržddlm} i }|j
                  j                  D ](  }t        ||«      sŒ|j                  ||j                  <   Œ* |j
                  j                  D ]<  }t        ||«      sŒ|j                  |v sŒ||j                     |j                  k7  sŒ< y y)NFr   )Ú	MemoryDepT)
rî  Úcan_fuser   Útorch._inductor.dependenciesr�  rû  Úreadsr§   r  r£  )r$   Únode1Únode2r�  Úreads1r  r  s         €r)   r‚  zPallasScheduling.can_fuse<  s¼   ø€ Ü‰wÑ  uÔ-Øð ×ÑÔ E×$6Ñ$6Ô$8Ý>àˆFØ×(Ñ(×.Ñ.ò 1�Ü˜c 9Õ-Ø'*§y¡y�F˜3Ÿ8™8Ò$ð1ð ×(Ñ(×.Ñ.ò %�Ü˜c 9Õ-°#·(±(¸fÒ2DØ˜cŸh™hÑ'¨3¯9©9Ó4Ù$ð%ð r+   c                óº  — t         j                  j                  }||j                  v r|j                  |   S t        j
                  j                  r$t        |t        j
                  j                  «      nd}t        j                  |j                  d«      «      j                  «       d d }|dk(  rd|› �}nd|› d|› �}||j                  |<   |j                  d|«      }t        «       }|j                  d|›d	�«       |j                  |d
¬«       |j                  d«       t!        ||«      \  }	}
|	› d|
› �}|j#                  ||j%                  «       |«       |S )NrR  zutf-8é   ÚfusedÚpallas_rE  ro  zasync_compile.pallas(z, r'''TrÚ  z''')rM  )r   r·   re  Úsrc_to_kernelr   ÚtritonÚdescriptive_namesr   ÚhashlibÚsha256ÚencodeÚ	hexdigestrý  r   rx  rÜ  r   Údefine_kernelr”  )r$   Úsrc_codeÚnode_schedulerÇ   rf  Ú
fused_nameÚkernel_hashrÕ  Úcompile_wrapperÚoriginsÚdetailed_originsÚmetadata_comments               r)   r“  zPallasScheduling.define_kernelT  sZ  € ô —'‘'×&Ñ&ˆØ�w×,Ñ,Ñ,Ø×(Ñ(¨Ñ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¸ÀwÓ$OÑ!ˆÐ!Ø%˜Y bÐ)9Ð(:Ð;ÐØ×Ñ˜k¨?×+CÑ+CÓ+EÐGWÔXàÐr+   )r  ztorch.devicer9   zOrderedSet[BackendFeature])r”  r:   r•  zSequence[BaseSchedulerNode]rÇ   rë  r9   r:   )r;   r<   r=   rë  Úkernel_typeÚclassmethodr  r‚  Úcan_fuse_verticalÚcan_fuse_horizontalr“  rx  ry  s   @r)   r{  r{  3  sZ   ø„ Ø€KàòEó ðEô
ð* !ÐØ"Ðð"àð"ð 3ð"ð ð	"ð
 
÷"r+   r{  )FÚ
__future__r   Údataclassesr�  r©   r[  rv  Útypingr   r   r   r  rµ   Útorch.utils._ordered_setr   Útorch.utils._sympy.functionsr   rƒ  r	   r‚  r
   rR  r   Úruntime.runtime_utilsr   rÆ   r   r   Úvirtualizedr   Úblock_analysisr   Úcommonr   r   r   r   r   Úsimdr   r   r   r   r"   ru  Úcollections.abcr@   rA   ÚirrB   Úops_handlerrC   Ú	schedulerrD   ÚMAIN_SUFFIXÚ_loggingÚgetArtifactLoggerr;   rL   rH   rŒ  rV   rX   Ú	dataclassr»  rÇ  rË  rÓ  rë  r{  r?   r+   r)   ú<module>r²     s`  ðÝ "ã Û Û Û 	Û ß +Ñ +ã ã Ý /÷ñ õ Ý 6ß >Ý Ý /÷õ ÷ CÑ Bô�Mô ñ< ‹×&Ñ&€ñ ß2åÝ+Ý-ð €ð —.‘.×2Ñ2°8¸]ÓK€÷>ñ >ô.U�,ô Uôq
˜Kô q
ðh ×Ñ÷ð ó ðð ×Ñ÷ð ó ðð ×Ñ÷ð ó ðð ×Ñ÷'ð 'ó ð'ô.W>'�:ô W>'ôt|C�~õ Cr+   