qi ddlmZddlZddlZddlmZmZmZmZddl Z ddl Z ddl m Z ddl mZddlmZddlmZdd lmZmZdd lmZd d lmZd d lmZmZmZmZd dlm Z m!Z!m"Z"erddl#m$Z$m%Z%ddl&m'Z'ddl(m)Z)ddl*m+Z+dZ,e jZj]e/dZ0GddZ1Gdde2Z3GddeZ4Gdde!Z5Gdde"Z6y)) annotationsN)AnyOptional TYPE_CHECKINGUnion) OrderedSet)has_tpu_pallas)configtorch_dtype_to_jax)get_fused_kernel_nameget_kernel_metadata)V)BlockPatternMatcher)BackendFeature CSEVariableIndentedBuffer OpOverrides)pexpr SIMDKernelSIMDScheduling)CallableSequence)IRNode) ReductionType)BaseSchedulerNodemain kernel_codec.eZdZdZ d ddZdddZy)PallasKernelWrapperz6Wrapper to provide .run() interface for Pallas kernelsNcL||_||_tjd|y)NzPallas kernel path: %s) kernel_fn kernel_pathkernel_code_loginfo)selfr$r%s d/home/ubuntu/crypto_trading_bot/.venv/lib/python3.12/site-packages/torch/_inductor/codegen/pallas.py__init__zPallasKernelWrapper.__init__(s$#&5{C)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 r,)r$)r(r,argskwargss r)runzPallasKernelWrapper.run/st~~t=F=f==r+N)r$zCallable[..., Any]r% Optional[str])__name__ __module__ __qualname____doc__r*r0r+r)r"r"%s.@KOD+D:GD!% >r+r"ceZdZdZy) UnsupportedzJException raised when an operation is not supported by the Pallas backend.N)r3r4r5r6r7r+r)r9r9>sTr+r9ceZdZdZedsdZedsdZedsdZedsdZedsdZ edsdZ edsdZ edsd Z edsd Z edsd Zedsd Zedsd ZedsdZedsdZedsdZedsdZedsdZedsdZedsdZedsdZedsdZedsdZedsdZedsdZedsdZedsdZedtdZedtdZ edtdZ!edudZ"e dv dwd!Z#edxd"Z$edyd#Z%edzd$Z&edsd%Z'edsd&Z(edsd'Z)edsd(Z*edsd)Z+edsd*Z,edtd+Z-edtd,Z.edtd-Z/edtd.Z0edtd/Z1edsd0Z2edsd1Z3edsd2Z4edtd3Z5edtd4Z6edtd5Z7edsd6Z8edtd7Z9edtd8Z:edtd9Z;edtd:Zedtd=Z?ed{d>Z@ed{d?ZAedsd@ZBedsdAZCedsdBZDedsdCZEedsdDZFedsdEZGedsdFZHedsdGZIedsdHZJedsdIZKedsdJZLedsdKZMedsdLZNedsdMZOedsdNZPedsdOZQed|dPZRed|dQZSed|dRZTed|dSZUed|dTZVedsdUZWed|dVZXed|dWZYed|dXZZed}dYZ[ed}dZZ\ed}d[Z]ed}d\Z^ed}d]Z_ed}d^Z`ed}d_Zaed}d`Zbed}daZced}dbZded}dcZeed}ddZfedsdeZgedsdfZhed~dgZiedtdhZjedtdiZkedtdjZledsdkZmedsdlZnedtdmZoedtdnZpedtdoZqedsdpZredtdqZsedtdrZty )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. cd|dS)Nzjnp.sin()r7xs r)sinzPallasKernelOverrides.sinJ!Ar+cd|dS)Nzjnp.cos(r=r7r>s r)coszPallasKernelOverrides.cosNrAr+cd|dS)Nzjnp.tan(r=r7r>s r)tanzPallasKernelOverrides.tanRrAr+cd|dS)Nz jnp.sinh(r=r7r>s r)sinhzPallasKernelOverrides.sinhV1#Qr+cd|dS)Nz jnp.cosh(r=r7r>s r)coshzPallasKernelOverrides.coshZrHr+cd|dS)Nz jnp.tanh(r=r7r>s r)tanhzPallasKernelOverrides.tanh^rHr+cd|dS)Nz jnp.arcsin(r=r7r>s r)asinzPallasKernelOverrides.asinbQCq!!r+cd|dS)Nz jnp.arccos(r=r7r>s r)acoszPallasKernelOverrides.acosfrOr+cd|dS)Nz jnp.arctan(r=r7r>s r)atanzPallasKernelOverrides.atanjrOr+cd|dS)Nzjnp.exp(r=r7r>s r)expzPallasKernelOverrides.expnrAr+cd|dS)Nz jnp.exp2(r=r7r>s r)exp2zPallasKernelOverrides.exp2rrHr+cd|dS)Nz jnp.expm1(r=r7r>s r)expm1zPallasKernelOverrides.expm1vA3a  r+cd|dS)Nzjnp.log(r=r7r>s r)logzPallasKernelOverrides.logzrAr+cd|dS)Nz jnp.log10(r=r7r>s r)log10zPallasKernelOverrides.log10~rZr+cd|dS)Nz jnp.log2(r=r7r>s r)log2zPallasKernelOverrides.log2rHr+cd|dS)Nz jnp.log1p(r=r7r>s r)log1pzPallasKernelOverrides.log1prZr+cd|dS)Nz jnp.sqrt(r=r7r>s r)sqrtzPallasKernelOverrides.sqrtrHr+cd|dS)Nz(1.0 / jnp.sqrt())r7r>s r)rsqrtzPallasKernelOverrides.rsqrts!!B''r+cd|dS)Nzjnp.abs(r=r7r>s r)abszPallasKernelOverrides.absrAr+cd|dS)Nz(-r=r7r>s r)negzPallasKernelOverrides.negsA3ayr+cd|dS)Nz jnp.floor(r=r7r>s r)floorzPallasKernelOverrides.floorrZr+cd|dS)Nz jnp.ceil(r=r7r>s r)ceilzPallasKernelOverrides.ceilrHr+cd|dS)Nz jnp.trunc(r=r7r>s r)trunczPallasKernelOverrides.truncrZr+cd|dS)Nz jnp.round(r=r7r>s r)roundzPallasKernelOverrides.roundrZr+cd|dS)Nz(1.0 / (1.0 + jnp.exp(-z)))r7r>s r)sigmoidzPallasKernelOverrides.sigmoids(3//r+cd|dS)N jnp.maximum(z, 0)r7r>s r)reluzPallasKernelOverrides.relusaS%%r+cd|d|dS)Nz jnp.power(, r=r7abs r)powzPallasKernelOverrides.powA3b1%%r+cd|d|dS)Nrwrzr=r7r{s r)maximumzPallasKernelOverrides.maximumaS1#Q''r+cd|d|dS)Nz jnp.minimum(rzr=r7r{s r)minimumzPallasKernelOverrides.minimumrr+cd|d|d|dS)N jnp.where(rzr=r7)condr|r}s r)wherezPallasKernelOverrides.wheresD6A3b1--r+Nc*t|}d|d|dS)Nz jnp.asarray( ).astype(r=r )r?dtype src_dtypeuse_compute_types jax_dtypes r)to_dtypezPallasKernelOverrides.to_dtypes#'u- aS )A66r+cFt|}t|}d|d|d|dS)z=Bitcast a value from one dtype to another with the same size.z)jax.lax.bitcast_convert_type(jnp.asarray(rz), r=r )r?rrr jax_src_dtypes r)to_dtype_bitcastz&PallasKernelOverrides.to_dtype_bitcasts7'u- *95 :1#Y}oUXYbXccdeer+cFddlm}tjj tjj |}tjj jtjj|||}tj||S)z>Convert a sympy expression to a JAX array indexing expression.r )get_bounds_index_expr)bounds) utilsrrkernelkexprprepare_indexingcsegeneratecomputer;r)exprrridx_strvars r) index_exprz PallasKernelOverrides.index_exprsq 2((..!:!:4!@Ahhll## HH  g.CD.I$ %--c599r+ct|}|tjk(r|rdSdSt|tr4t j |ryt j|r |dkDrdSdSd|d|d S) z/Convert a constant value to JAX representation.TrueFalsezjnp.nanrzjnp.infz-jnp.infz jnp.array(z, dtype=r=)r torchbool isinstancefloatmathisnanisinf)valrrs r)constantzPallasKernelOverrides.constantsp'u- EJJ  6 -g - c5 !zz# zz#$'!Gy;;C5 155r+cd|dS)Nz jnp.real(r=r7r>s r)realzPallasKernelOverrides.realrHr+cd|dS)Nz jnp.imag(r=r7r>s r)imagzPallasKernelOverrides.imagrHr+cd|dS)Nz jnp.conj(r=r7r>s r)conjzPallasKernelOverrides.conjrHr+cd|dS)Nz jnp.angle(r=r7r>s r)anglezPallasKernelOverrides.anglerZr+cd|d|dS)z8View complex tensor as real tensor with extra dimension.zjnp.stack([jnp.real(z ), jnp.imag(z )], axis=-1)r7r>s r) view_as_realz"PallasKernelOverrides.view_as_reals&aS QC|DDr+cd|d|dS)z#View real tensor as complex tensor.(z[..., 0] + 1j * z [..., 1])r7r>s r)view_as_complexz%PallasKernelOverrides.view_as_complexs1#%aS 22r+cd|d|dS)Nrz == r=r7r{s r)eqzPallasKernelOverrides.eq 1#T!Ar+cd|d|dS)Nrz != r=r7r{s r)nezPallasKernelOverrides.ne rr+cd|d|dS)Nrz < r=r7r{s r)ltzPallasKernelOverrides.lt1#S1~r+cd|d|dS)Nrz <= r=r7r{s r)lezPallasKernelOverrides.lerr+cd|d|dS)Nrz > r=r7r{s r)gtzPallasKernelOverrides.gtrr+cd|dS)Nz jnp.isnan(r=r7r>s r)rzPallasKernelOverrides.isnanrZr+cd|dS)Nz jnp.isinf(r=r7r>s r)rzPallasKernelOverrides.isinf!rZr+cd|dS)Nz jnp.isfinite(r=r7r>s r)isfinitezPallasKernelOverrides.isfinite%sqc##r+cd|d|dS)Nrz >= r=r7r{s r)gezPallasKernelOverrides.ge)rr+cd|d|dS)Nzjnp.logical_and(rzr=r7r{s r) logical_andz!PallasKernelOverrides.logical_and.!!Bqc++r+cd|d|dS)Nzjnp.logical_or(rzr=r7r{s r) logical_orz PallasKernelOverrides.logical_or2 2aS**r+cd|dS)Nzjnp.logical_not(r=r7r>s r) logical_notz!PallasKernelOverrides.logical_not6!!A&&r+cd|d|dS)Nzjnp.logical_xor(rzr=r7r{s r) logical_xorz!PallasKernelOverrides.logical_xor:rr+cd|d|dS)Nz jnp.arctan2(rzr=r7r{s r)atan2zPallasKernelOverrides.atan2?rr+cd|d|dS)Nz jnp.hypot(rzr=r7r{s r)hypotzPallasKernelOverrides.hypotCrr+cd|d|dS)Nz jnp.fmod(rzr=r7r{s r)fmodzPallasKernelOverrides.fmodGs1#Rs!$$r+cd|d|dS)Nzjnp.remainder(rzr=r7r{s r) remainderzPallasKernelOverrides.remainderKs"QCq))r+c &d|d|d|d|d|d S)Nz (jnp.sign(z ) * jnp.sign(z ) * (jnp.abs(z ) // jnp.abs(z ))).astype(.dtype)r7r{s r)truncdivzPallasKernelOverrides.truncdivOs.A3mA3mA3mA3kZ[Y\\cddr+cd|d|dS)Nrz // r=r7r{s r)floordivzPallasKernelOverrides.floordivUrr+cd|d|d|dSNz jnp.clip(rzr=r7r?min_valmax_vals r)clampzPallasKernelOverrides.clampY1#Ry7)155r+cd|d|d|dSrr7rs r)clipzPallasKernelOverrides.clip]rr+cd|dS)Nz jnp.sign(r=r7r>s r)signzPallasKernelOverrides.signbrHr+cd|dS)Nz jnp.signbit(r=r7r>s r)signbitzPallasKernelOverrides.signbitfsaS""r+cd|dS)Nzjax.scipy.special.erf(r=r7r>s r)erfzPallasKernelOverrides.erfks's!,,r+cd|dS)Nzjax.scipy.special.erfc(r=r7r>s r)erfczPallasKernelOverrides.erfcos(1--r+cd|dS)Nzjax.scipy.special.erfinv(r=r7r>s r)erfinvzPallasKernelOverrides.erfinvss*1#Q//r+cd|dS)Nzjax.scipy.special.gammaln(r=r7r>s r)lgammazPallasKernelOverrides.lgammaw+A3a00r+cd|dS)Nzjax.scipy.special.digamma(r=r7r>s r)digammazPallasKernelOverrides.digamma{rr+cd|d|d|dS)Nrz>.astype(jnp.float64) == 0.0, 1.0, jax.scipy.special.bessel_jn(z&.astype(jnp.float64), v=0)[0]).astype(rr7r>s r) bessel_j0zPallasKernelOverrides.bessel_j0+++,#.c " r+cd|d|d|dS)Nrz>.astype(jnp.float64) == 0.0, 0.0, jax.scipy.special.bessel_jn(z&.astype(jnp.float64), v=1)[1]).astype(rr7r>s r) bessel_j1zPallasKernelOverrides.bessel_j1rr+cd|d|dSNjax.lax.bessel_i0e() * jnp.exp(jnp.abs(rfr7r>s r)modified_bessel_i0z(PallasKernelOverrides.modified_bessel_i0%QC';A3bAAr+cd|d|dSNjax.lax.bessel_i1e(rrfr7r>s r)modified_bessel_i1z(PallasKernelOverrides.modified_bessel_i1rr+cd|d|d|dS)Nrz == 0.0, 1.0, jnp.sin(z) / r=r7r>s r)spherical_bessel_j0z)PallasKernelOverrides.spherical_bessel_j0s A34QCtA3a@@r+cd|d|dSrr7r>s r)i0zPallasKernelOverrides.i0%QC';A3bAAr+cd|dS)Nrr=r7r>s r)i0ezPallasKernelOverrides.i0e%QCq))r+cd|d|dSr r7r>s r)i1zPallasKernelOverrides.i1rr+cd|dS)Nr r=r7r>s r)i1ezPallasKernelOverrides.i1err+cd|d|dSNzjax.scipy.special.gammainc(rzr=r7r?ys r)gammainczPallasKernelOverrides.gammaincs-QCr!A66r+cd|d|dSNzjax.scipy.special.gammaincc(rzr=r7rs r) gammaincczPallasKernelOverrides.gammaincc.aS1#Q77r+cd|d|dSrr7rs r)igammazPallasKernelOverrides.igammas-QCr!A66r+cd|d|dSrr7rs r)igammaczPallasKernelOverrides.igammacr r+cd|d|dS)Nzjax.scipy.special.polygamma(z.astype(jnp.int32), r=r7rs r) polygammazPallasKernelOverrides.polygammas.aS0DQCqIIr+cd|dS)Nzjax.scipy.special.ndtri(r=r7r>s r)ndtrizPallasKernelOverrides.ndtris*!A..r+cd|d|dS)Nzjax.scipy.special.zeta(rzr=r7rs r)zetazPallasKernelOverrides.zetas)2aS22r+cd|d|dS)Nzjax.scipy.special.xlogy(rzr=r7rs r)xlogyzPallasKernelOverrides.xlogys*!Bqc33r+cd|d|dS)Nzjax.scipy.special.xlog1py(rzr=r7rs r)xlog1pyzPallasKernelOverrides.xlog1pys,A3b155r+c>d|d|d|d|d|d|d|d|d |d S) Njnp.where(jnp.abs() <= 1, jnp.cos( * jnp.arccos(jnp.clip(, -1, 1))), jnp.where( > 1, jnp.cosh( * jnp.arccosh(jnp.maximum(, 1.0))), ((-1.0) ** ) * jnp.cosh( * jnp.arccosh(jnp.maximum(- , 1.0)))))r7r?ns r)chebyshev_polynomial_tz,PallasKernelOverrides.chebyshev_polynomial_ts[!$c04s5aS9M!,H:  W r+cdjgd|d|d|d|d|d|d|d |d |d |d |d |d|d|d|d|d |dS)Nr0) < 1, jnp.sin(( + 1) * jnp.arccos(jnp.clip(z&, -1, 1))) / jnp.sqrt(jnp.maximum(1 - z**2, 1e-10)), jnp.where( >= 1, jnp.where( == 1,  + 1.0, jnp.sinh(( + 1) * jnp.arccosh(jnp.maximum( , 1.0))) / jnp.sqrt(jnp.maximum(**2 - 1, 1e-10))), jnp.where( == -1, ((-1.0) ** ) * ( + 1.0), ((-1.0) ** ) * jnp.sinh((! + 1) * jnp.arccosh(jnp.maximum(-**2 - 1, 1e-10)))))joinr:s r)chebyshev_polynomial_uz,PallasKernelOverrides.chebyshev_polynomial_us ; ; ; ;$ ;s ;6 ;78c ;:( ;()s ;+ ; ; ; ;# ;$%# ;& ;  ; < ; =>3 ; ?$ ; %&3 ; ' ; ;/ ;01c ;27 ;89c ;: ; ;+ ;,-# ;.O ;PQc ;R$ ;%&3 ;': ; r+cdjgd|d|d|d|d|d|d|d|d |d |d |d |d |d|d|d|d |d|d|d|d|d|d|dS)Nr>r == 0, jnp.ones_like( ), jnp.where( == 1, 2* - 1, jnp.where( == 2, 4***2 - 2* == 3, 8***3 - 4***2 - 4* + 1, jnp.where( == 4, 16***4 - 8* **3 - 12***2 + 4* == 5, 32* **5 - 16* **4 - 32* **3 + 12***2 + 6* - 1, jnp.zeros_like()))))))rMr:s r)chebyshev_polynomial_vz,PallasKernelOverrides.chebyshev_polynomial_v  ) )j ) )0 ) )4 ) )$ )%&C )( ) )$ )%&C )'/ )01s )3 ) )% )&'C )(0 )12s )3; )<=# )> )  )&  )'(S  ))1  )23  )4=  )>?C  )@H  )IJs  )K  )   ) &  ) '(S  ) )2  ) 34  ) 5>  ) ?@S  ) AJ  ) KL  ) MU  ) VWTW  ) X  )  S  ) !(  ) r+cdjgd|d|d|d|d|d|d|d |d |d |d |d |d |d|d|d |d|d|d|d|d|d|d|dS)Nr>rrQrRrSrZrU**2 + 2*rTrW**3 + 4*rYr[**4 + 8*r]r_ **5 + 16*rarc + 1, jnp.zeros_like(rerMr:s r)chebyshev_polynomial_wz,PallasKernelOverrides.chebyshev_polynomial_wrgr+cJd|d}d|d|d|d|d|d|d |d |d |d S) N(2 *  - 1)r0r1r2r3r4r5r6r7r8r9r7r?r;rs r)shifted_chebyshev_polynomial_tz4PallasKernelOverrides.shifted_chebyshev_polynomial_t/si A3e  $c04s5aS9M!,H:  W r+cd|d}djgd|d|d|d|d|d |d |d |d |d |d|d|d|d|d|d|d |dS)Nrprqr>r0r?r@z', -1, 1))) / jnp.sqrt(jnp.maximum(1 - (z)**2, 1e-10)), jnp.where(rArBrCrDrErFrGrHrIrJrKrLrMrrs r)shifted_chebyshev_polynomial_uz4PallasKernelOverrides.shifted_chebyshev_polynomial_u=sA3e  ; ; ; ;$ ;s ;6 ;78c ;:) ;)* ;, ; ; ; ;# ;$%# ;& ;  ; < ; =>3 ; ?$ ; %&3 ; ' ; ;/ ;01c ;27 ;89c ;: ; ;+ ;,-# ;.O ;PQc ;R$ ;%&3 ;': ; r+cd|d}djgd|d|d|d|d|d |d |d|d |d |d |d|d|d|d|d|d|d|d|d|d|d|d|dS)Nrprqr>rrQrRrSrTrUrVrWrXrYrZr[r\r]r^r_r`rarbrcrdrerMrrs r)shifted_chebyshev_polynomial_vz4PallasKernelOverrides.shifted_chebyshev_polynomial_vOA3e  ) )j ) )0 ) )4 ) )$ )%&C )( ) )$ )%&C )'/ )01s )3 ) )% )&'C )(0 )12s )3; )<=# )> )  )&  )'(S  ))1  )23  )4=  )>?C  )@H  )IJs  )K  )   ) &  ) '(S  ) )2  ) 34  ) 5>  ) ?@S  ) AJ  ) KL  ) MU  ) VWTW  ) X  )  S  ) !(  ) r+cd|d}djgd|d|d|d|d|d |d |d |d |d |d|d |d|d|d|d|d|d|d|d|d|d|d|dS)Nrprqr>rrQrRrSrZrUrirTrWrjrYr[rkr]r_rlrarcrmrerMrrs r)shifted_chebyshev_polynomial_wz4PallasKernelOverrides.shifted_chebyshev_polynomial_w]rxr+cdjgd|d|d|d|d|d|d|d |d |d|d |d |d |d|d|d|d|dS)Nr>rrQrRz == 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(rerMr:s r)hermite_polynomial_hz*PallasKernelOverrides.hermite_polynomial_hks~ ) )j ) )0 ) )4 ) )& )'(c )* ) )& )'(c )* ) )' )()c )*5 )67C )8 )  )(  ))*s  )+6  )78S  )9  )   ) (  ) )*s  ) +7  ) 89c  ) :F  ) GHS  ) I  )  S  ) !(  ) r+cdjgd|d|d|d|d|d|d|d |d |d|d |d |d |d|d|d|d|dS)Nr>rrQrRrBr|z == 2, z**2 - 1, jnp.where(z == 3, **3 - 3 * z == 4, z **4 - 6 * z**2 + 3, jnp.where(z == 5, z **5 - 10 * **3 + 15 * r}rerMr:s r)hermite_polynomial_hez+PallasKernelOverrides.hermite_polynomial_he}s}  ) )j ) )0 ) )4 ) )" )#$# )& ) )" )#$# )& ) )# )$%# )&0 )12s )3 )  )#  )$%#  )&0  )12s  )3  )   ) #  ) $%#  ) &1  ) 23  ) 4?  ) @Ac  ) B  )  S  ) !(  ) r+cdjgd|d|d|d|d|d|d|d |d |d |d |d |d|d|d|d|d|d|d|d|d|d|d|dS)Nr>rrQrRz == 1, 1 - r|z == 2, (rYz + 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(rerMr:s r)laguerre_polynomial_lz+PallasKernelOverrides.laguerre_polynomial_ls   ) )j ) )0 ) )4 ) )& )'(c )* ) )# )$%3 )&. )/0c )2 ) )% )&'C )(0 )12s )3< )=>3 )? )  )$  )%&3  )'0  )12s  )3<  )=>3  )?H  )IJs  )K  )   ) %  ) &'C  ) (1  ) 23  ) 4>  ) ?@S  ) AK  ) LM#  ) NX  ) YZWZ  ) [  )  S  ) !(  ) r+cdjgd|d|d|d|d|d|d|d |d |d |d |d |d|d|d|d|d|dS)Nr>rrQrRrBr|z == 2, (3 * z**2 - 1) / 2, jnp.where(z == 3, (5 * rz) / 2, jnp.where(z == 4, (35 * z **4 - 30 * z**2 + 3) / 8, jnp.where(z == 5, (63 * z **5 - 70 * rz) / 8, jnp.zeros_like(rerMr:s r)legendre_polynomial_pz+PallasKernelOverrides.legendre_polynomial_ps~  ) )j ) )0 ) )4 ) )" )#$# )& ) )' )()s )+ ) )( ))*s )+5 )67C )8 )  ))  )*+  ),7  )89c  ):  )   ) )  ) *+  ) ,7  ) 89c  ) :E  ) FGC  ) H  )  S  ) !(  ) r+cd|dS)Nzjnp.reciprocal(r=r7r>s r) reciprocalz PallasKernelOverrides.reciprocals 1%%r+cd|dS)Nz jnp.square(r=r7r>s r)squarezPallasKernelOverrides.squarerOr+cd|d|d|dS)zFused multiply-add: a * b + czjnp.fma(rzr=r7)r|r}cs r)fmazPallasKernelOverrides.fmas!BqcA3a((r+cd|d|dS)Nz jnp.copysign(rzr=r7r{s r)copysignzPallasKernelOverrides.copysignsqcA3a((r+cd|d|dS)Nzjnp.nextafter(rzr=r7r{s r) nextafterzPallasKernelOverrides.nextafterrr+cd|d|dS)Nz jnp.ldexp(rzr=r7r{s r)ldexpzPallasKernelOverrides.ldexprr+cd|dS)Nz jnp.frexp(r=r7r>s r)frexpzPallasKernelOverrides.frexprZr+cd|dS)Nz jnp.modf(r=r7r>s r)modfzPallasKernelOverrides.modfrHr+cd|d|dS)Nzjnp.bitwise_and(rzr=r7r{s r) bitwise_andz!PallasKernelOverrides.bitwise_andrr+cd|d|dS)Nzjnp.bitwise_or(rzr=r7r{s r) bitwise_orz PallasKernelOverrides.bitwise_orrr+cd|d|dS)Nzjnp.bitwise_xor(rzr=r7r{s r) bitwise_xorz!PallasKernelOverrides.bitwise_xorrr+cd|dS)Nzjnp.bitwise_not(r=r7r>s r) bitwise_notz!PallasKernelOverrides.bitwise_notrr+cd|d|dS)Nzjnp.left_shift(rzr=r7r{s r) left_shiftz PallasKernelOverrides.left_shiftrr+cd|d|dS)Nzjnp.right_shift(rzr=r7r{s r) right_shiftz!PallasKernelOverrides.right_shiftrr+)r?strreturnr)r|rr}rrr)rrr|rr}rrr)NT) r?rr torch.dtyperzOptional[torch.dtype]rrrr)r?rrrrrrr)r sympy.Exprrrrr)rrrr)r?rrrrrrr)r?rrrrr)r?rr;rrr)r|rr}rrrrr)ur3r4r5r6 staticmethodr@rCrErGrJrLrNrQrSrUrWrYr\r^r`rbrdrgrirkrmrorqrsrurxr~rrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrr r rrrrrrr"r$r&r(r*r,r.r<rOrfrnrsrurwrzr~rrrrrrrrrrrrrrrrrr7r+r)r;r;Bs       """"""  !!!!  !!  ((!!  !!!!00&&&&((((..,0"& 7 77)7 7 77ff:: 6 6      !!EE33 !!!!$$,,++'',,((&&%%**ee 6666  ##--..001111        BB BB AA BB**BB**77 887788JJ //334466      (              "          "            &&""))))**&&!!  ,,++,,''++,,r+r;ceZdZUdZeZeZded<fdZ ddZ ddZ ddZ ddZ dd Zdd Zdd Zdd Zdd ZddZddZ d ddZ ddZeddZdd dZdd!dZxZS)" PallasKernelad Pallas kernel for elementwise operations with support for strided/scatter access. Strategy: - Convert index expressions to JAX-compatible array slicing - Load/store using indexed access: "in_ptrX[slice]" or full-array "in_ptrX[...]" - Compute expression with Python operators (compatible with jax.numpy broadcasting) - Generate Python code that defines a Pallas kernel and a host entrypoint. - Use async_compile.pallas path to compile and load Python code. For GPU (Triton backend): - Use masked loads/stores with power-of-2 block sizes to handle non-power-of-2 shapes zCallable[[sympy.Expr], str]rct||i|tjj }|j dk(|_d|_i|_g|_ i|_ y)Ncuda) superr*rgraphget_current_device_or_throwtypeis_gpuuse_masked_ops tensor_masksstore_with_outputload_index_exprs)r(r.r/device __class__s r)r*zPallasKernel.__init__sW $)&)446kkV+ +/8:79r+cy)z)Check array bounds for indirect indexing.Nr7)r(rsizeloweruppers r) check_boundszPallasKernel.check_boundssr+c|j|}t|tjry|jr 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 ...)rrsympySymbol is_Integerr_convert_to_jax_slice)r(indexprepared_indexs r)_get_index_strzPallasKernel._get_index_str sN"..u5 nell 3  & &~& &--n= =r+c|jsytjjj |}|j }t |jj}||z}t|dk(r t|St|dk(rtt|}tj||}tj||}|}||z }tjjj |}|dk(r|dk(ry|dk(r|j!|} d| S|j!|} |j!|} | d| S||z }tjjj |}|dk(rV||k(rQyt|dkDrBd} |D]7}tj||}tj||}|dk7s5d} n| ryy|j#|S)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. rrrz::TF) range_treesrrsizevarssimplify free_symbolsrrange_tree_nodeskeyslenrnextiterrget_subexpr_involving_symbolmatch_affine_block_exprr_generate_index_array) r(rr iter_vars used_varsrvar_exprstrideoffset stride_str offset_strall_unit_strides r)rz"PallasKernel._convert_to_jax_slice(s  ))%0)) t4499;< !9, y>Q u:  ^q tI'C+GGsSH)@@3OF!)))226:Q;6Q; q[!%F!3J |,,"&F!3J!%F!3J(\J<88)))226:Q;8s? ^a #O  .KKESVW,DDXsSQ;&+O   ))%00r+ctd|)z Generate JAX code to compute an index array for complex indexing patterns. For very complex patterns that can't be expressed as simple slices, we need to compute the indices explicitly. This is not yet fully implemented. z>Pallas backend does not yet support complex indexing pattern: )r9)r(rs r)rz"PallasKernel._generate_index_array~sLUG T  r+c||j}t|jj}t ||zS)zFCheck if index expression contains iteration variables (x0, x1, etc.).)rrrrr)r(rrrs r)_has_iteration_varsz PallasKernel._has_iteration_varss6)) t4499;< L9,--r+cb|j}|D]}t|jdsyy)zICheck if index expression contains indirect variables (tmp0, tmp1, etc.).tmpTF)rr startswith)r(rrsyms r)_has_indirect_varszPallasKernel._has_indirect_varss6))  C3x""5) r+c|j|}|j|}|r|r|j|dfS|r|j|dfS|j |dfS)a Get the index expression string and whether it needs flattening. Returns: Tuple of (index_str, needs_flatten) where needs_flatten indicates if the buffer should be flattened before indexing (for mixed indexing). TF)rr_handle_mixed_indexingrr)r(r has_indirect has_iter_varss r)_get_index_exprzPallasKernel._get_index_exprsl..u5 007 M..u5t; ; ::e$e+ +&&u-u4 4r+c" |jsyt}|j|jjj |j|jj j |jtjjj g}|D]} tjj|}|j}td|D}d}|D]"}t|dr|t|z}||z}$|j|||f|rAt#|dkDr3t%d|D} | sy|dd t% fd|D} | Sy#t $rYwxYw) aB Determine if we should use masked ops for this entire kernel. Masked ops with pl.ds(block_size) flatten tensors to 1D, which works when: 1. We're on GPU (CUDA backend uses Triton which requires power-of-2 sizes) 2. All tensors are already 1D (so flattening doesn't change dimensionality) 3. All tensors have the same size (so broadcasting works correctly) With per-tensor masks, each tensor gets its own mask based on its size. This should be called once in codegen_kernel() before generating the kernel body. Fc3NK|]}t|dr t|n|yw__int__Nhasattrint.0ss r) z@PallasKernel._determine_masked_ops_for_kernel..s"S9(=c!f1DS#%rrrc3@K|]\}}}t|dk(ywrN)r)r_shapes r)rz@PallasKernel._determine_masked_ops_for_kernel..sE[QqUqEsr c30K|] \}}}|k(ywr1r7)rrr first_sizes r)rz@PallasKernel._determine_masked_ops_for_kernel..sNzq!T 2Ns)rrupdater. input_buffersroutput_buffersrrname_to_buffer get_bufferget_sizetuplerrappend Exceptionrall) r(all_buffer_namesbuf_infobuf_namebufrr total_sizerall_1d all_same_sizers @r) _determine_masked_ops_for_kernelz-PallasKernel._determine_masked_ops_for_kernelst{{&<  7 7 < < >? 8 8 = = ?@ 6 6 ; ; =>( H gg((2||~SdSS (Aq),"c!f, "a ( 5* => * H )EHEEF"!QJNXNNM '  s>A>F F Fcd||jvrd|}||j|<|j|S)z2Get or create a unique mask variable for a buffer.mask_)r)r(rmask_vars r)_get_or_create_maskz PallasKernel._get_or_create_masks= 4,, ,xj)H*2D  h '  **r+c|jj|}tjj |}||j |<|j |j|_|j|\}}|dk(xr| xr |j }|r|j|}d|d|d} n|r |d|d} n|d|d} |jj|j| |S) Nrzpltriton.load(z.at[pl.ds(block_size)], mask=r=z[...].flatten()[][r) r.inputrr get_dtyperrr rrrrr) r(namerrr index_str needs_flatten use_maskedr load_exprs r)loadzPallasKernel.loads iiood#!!$'',d#    &"&"G"G"ID #'#7#7#> =%'U ,=U$BUBU //5H(-J8*TUVI %/ {!Extract the coefficient of a variable in the index expression.r)coeffrdiffr TypeError ValueError)rrrs r)get_coefficientz.get_coefficient'sNKK$Ez 5#. 5z!z*  s :A  A T)keyreverserrindirectc |dSNrr7r>s r)z5PallasKernel._handle_mixed_indexing..Ds !A$r+z jnp.arange(r=c3.K|] }|ks dywrr7rr var_coeffs r)rz6PallasKernel._handle_mixed_indexing..[s%NAI a%N c3.K|] }|ks dyw)r Nr7r*s r)rz6PallasKernel._handle_mixed_indexing..\s*1y=A*r,z, Nonez[:rc3.K|] }|kDs dywrr7rrindirect_coeffs r)rz6PallasKernel._handle_mixed_indexing..lsI!a.6HAIr,c3.K|] }|ks dywrr7r/s r)rz6PallasKernel._handle_mixed_indexing..msJ1q>7IQJr,zNone, rrz...]z[...)rrrrrrsortedrrrsort enumeratelengthsumvaluesreplace) r(rrrused_iter_vars_setr"used_iter_varsr iter_coeffsrrindirect_var_symsr indirect_varsindirect_coeffsall_componentsivar_name range_entry range_size arange_exprn_trailing_itern_trailing_indirect n_trailing trailing_dims indirect_var n_leading leading_nonestrailing_nones reshape_exprr0r+s ` @@r)rz#PallasKernel._handle_mixed_indexings<t4499;< )) )I5 ! "a '::e$ $  2QUV7EFs+F FJJu% (4Q1A8I8I%8PQQQ->?cS? ?@QQ!3q6?1#55QQ! GC  ! !?3#7"E F G$ KC  ! !?3#7S"I J K= / EFAs3xHd+++"33C8 (// +C0 +DJJz,B+C1E #&%N%N"N&)*.557*'#-/BB >$,z$9M%0MM?!"DK%--h D 1 E6* FL,\:NI{IIIJ JJJ1}a (9 4 !)J!6".qs>BRRST Q (9 4 ".qtD a!)J!6".tN3C1E + !)), EI+ F.WGR?Rs.JJ"4J">J'J,c| td|jj|}|jj ||j |j |_ tjj|}|j}t|dk(}|r|d|} ng|j|\} } | dk(xr| xr |j } | r|j|} d|d|d | d } n| dk(r |d|d |d } n |d | d|} |jj!| |j"j%|| fy#t$rd}d}YwxYw)Nzpallas store mode not supportedrr7Fz[...] = rzpltriton.store(z.at[pl.ds(block_size)], z, mask=r= .reshape(.shape)rz] = )r9r.outputstore_buffer_namesaddrr rrrrrrrrstores writelinerr)r(rrvaluemodeoutr output_shape is_scalar store_exprrrrrs r)storezPallasKernel.stores  ?@ @iit$ ##D)    &"&"G"G"ID  ''$$T*C<<>LL)Q.I 50J'+';';E'B $I}U"P='8PT=P=P 33D9.se3KE7RYZbYccde e#!$uHUG9SEI !$uAi[UG<  j) %%sJ&78E LI s&=E EEc jsJt|tr td|||f}|jj vrjj |Sdddddddd }t gd }tfd |D}g} tjjd D]F\} } | jjdr" | jt| jHd} |D]/} | j$vsj$| } | t|z} 1d}j$D]2} | jdsj$| } |t|z}4t'djjD}|dk(r|r| r |r d|d| d}nd|d}n|dvr}||}|xr | xr| dkD}|r_|dkDrYd}j(rt+t-j(j/}jjD cgc]"\} } | jjdr| $}} } |r|d}|j1|} |dk7r t|nd}jjD cgc]"\} } | jjds| $}} } |r2|d}|j1|} |dk7r t|nd}||kDrdnd}|dk(r |d|d|d}ntt3|}|d|d|d}np|d|d}ng||vr:|r,| r*|r(| r&t5| }||}d|d|d|d| d|d |d }n5||d|d}n)td!|d"t7|j9d#jj;j<||$} | jj |<| S#t t"f$rd} YwxYw#t t"f$rd} YwxYw#t t"f$rd}YwxYwcc} } w#t t"f$rd}YwxYwcc} } w#t t"f$rd}YpwxYw)%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. zHTuple reductions (e.g., welford_combine) not supported in Pallas backendzjnp.sumzjnp.prodzjnp.maxzjnp.minzjnp.anyz jnp.argmaxz jnp.argmin)r6prodmaxminanyargmaxargmin)r?rzc3:K|]}|jvywr1)numels)rpr(s r)rz)PallasKernel.reduction..sIA,Isct|dSr'rr>s r)r(z(PallasKernel.reduction..sQqTr+)r#rNrc3\K|]$\}}|jjdrd&yw)rjrN)prefixr)rrentrys r)rz)PallasKernel.reduction..s. U||&&s+  s*,xor_sumzjnp.bitwise_xor.reduce(rOz, -1), axis=-1)r=)rbrcrrz, axis=z(lambda v: (lambda pw_sizes: z(v.reshape(-1, z>), axis=-1) if v.ndim == 2 else (lambda input_shape, pw_axes: z=(jnp.moveaxis(v, pw_axes, list(range(len(pw_axes)))).reshape(z`, -1), axis=-1))(v.shape, [i for i, s in enumerate(v.shape) if s in pw_sizes][:len(pw_sizes)]))(z))(zReduction type 'z8' not yet supported in Pallas backend. Supported types: z , xor_sumr)inside_reductionrrr9rreduction_cacherrar2ritemsrlrrrr5r r!rfr6rrrr7rrangerlistrrr)!r(rrreduction_typerV cache_key reduction_opspointwise_prefixes has_pointwisepointwise_sizesrrmpointwise_numelrgnumelreduction_numeln_reduction_dimsreduction_expr reduction_opis_partial_reductionreduction_axis load_indexreduction_varsr_varr_coeffr_stridepw_varspw_varpw_coeff pw_strideaxes pw_sizes_strresults!` r) reductionzPallasKernel.reductionsP"$$$$ eU #Z  6 00 088++I6 6"" (8I6HII   ! ! ' ' )/B  JC<<**3/#**3u||+<=  # ADKK A#s5z1O   A||C  A#s5z1O   "3399;   Y &_#:5'?J[[j!k#:5'!C 3 3)8LI/Io6I !$(81(<"#((!%d4+@+@+G+G+I&J!KJ+/*?*?*E*E*G&&C <<2237&N& & .q 1","2"25"9)7>!|s7|H /3.C.C.I.I.K# *U#(<<#:#:3#? ## #%,QZF'1'7'7'?H.=E]CMPQ 3;Y2FQAN#q((4~QugW^DTTU%VN!'7!89D(4~QugWTF!%LN%1>5'!; } ,### #?3 ,^< 3#nOO3DE5$n%SSbRcd &c% 3%2.$A#B!E7!!L">"23$$(););)=$>#?yJ  "" LL # /5  + _":.&*O":.&*O":.&*OD&!*:6)'(H)#%.z#:.,- .sf$O*#P&P6'P59P;('Q+Q*P?PPPP21P2;QQQ.-Q.ctjj|}|j}|j Sr1)rrr get_layout is_contiguous) buffer_namerlayouts r)_buffer_is_contiguousz"PallasKernel._buffer_is_contiguous}s1gg  -!##%%r+c At}|jj\}}}}|Dcgc]}|j}}|Dcgc]}|j ds|}}|Dcgc]}|j ds|} }| s t d|jj jD cic]\} } t| tr| | } } } |xsd} tjjjdk(}tjj j"}|rDtjj j$s t dt's t d|rdnd }d |j(rd nd zd z|sdnd z|j(rdnd z}|j+|di}|D]4}| j-|}|duxr|j/|}| xr|||<6|Dcgc] }||s |d}}|Dcgc]}|j ds|}}||z}||z}t1|jDcgc] \}}|r | c}}}t3| Dcgc] \}}||vs |}}}||_d| ddj7||j(rdnd zdz}|j9||j;5|j(r6|j<r)i}|jj>jD]\} } t| tr| n| || <|jj jD]\} } t| tr| n| || <tA|j<jD]{\} }!|j-| | }"d}#|D]}|"|k(s| t|vs|}#n|#s:|j9d| |j9|!d|#d|j9|!d|!d}|jBrS|j(sF|j9dd}$d}%| ri| d }&| j-|&}'|'rQ tjjE|'}(|(jG})tId!|)D}$d"}%|$D]}*|%|*z}%  |jBjD]\}+},t|+}-|,jL}.|jO|.}/ tQ|.d#r tS|.nd}0|0K|$rFtY|$d"kDr8|0|%k(r3dj7d$|$D}1|j9|-d%|/d&|1d'|j9|-d%|/d'|jZj\D]}2|j9t|2|j^D]\}3}4|3|vs |j9|4 ddd| d(}5g}6t3|D]/\}}||vs|j d)s|6ja|d*z1|6rd+dj7d,|6Dzd-z}7nd.}7|j9d/|7d'|j9d|5d0dj7|d|j;5|j9d1|j9d2|j9d3|j9d'|j(r|j9d4|j9d5|j9d6|D]}|j9d7|d8|j9d9|j9d:|j9d;|j9d<g}8t3| D]w\}9}|j dr=|j-|d=s*|d}:|jc|:};|8ja|;|9fT|jc|};|8ja|;|9fydj7d>|8D}<|j(rd?| d@n| dA}=|j9dB|j9dC|=z|j9dD|j9dE|dF|j9dG|j9|8rdH|<dIndJ|j9dK|r$|j9dCdj7|dF|j9d'ddd| dL}>|j9d|>d+dj7|dM|j;5|j9dN|j9dO|rI|j9dP|D]3}:|r|j9|:dQ|:dR|j9|:dS|:dT5|j9dU|D]E}?|?j d)s|r|j9|?dQ|?dR/|j9|?dS|?dTG|j9dV|D]E}?|?j dWs|r|j9|?dQ|?dR/|j9|?dS|?dXG|j9dY|j9dZdj7| Dcgc]}d[|d\ c}zd-z|j9d]dj7| Dcgc]}d^|d_ c}zd-ziA|D] }:|:d`A|:< |D] }?|?d`A|?< |r3dj7Afda|D}@|j9db|5d0|@d'n|j9db|5dc|r`|j9dd|D]J}| |}|r*|j9de|df|j9|dg4|j9|dh|diLddd|jeScc}wcc}wcc}wcc} } wcc}wcc}wcc}}wcc}}w#tJ$rYJAX conversions via DLPack Args: name: Optional kernel name (will use placeholder if not provided) Returns: str: Complete Python source code for the Pallas kernel out_ptr)r in_out_ptrz2Pallas backend requires at least one output buffer cpuzBPallas backend currently only supports using the first JAX device.zPALLAS_TARGET_TPU is set, but no TPU device was found. Please make sure that you have a TPU available and that JAX is configured correctly.rrz* import functools zimport math r>zimport torch import jax import jax.numpy as jnp from jax.experimental import pallas as pl from torch._inductor.runtime.runtime_utils import torch_dtype_to_jax_runtime zC from jax.experimental.pallas import triton as pltritonzN from torch._inductor.runtime.runtime_utils import next_power_of_2TstripN_alias)rin_ptrzdef z_kernel(rzz, *, block_sizez):z # Mask for z_size = z.sizez = jnp.arange(block_size) < _sizez*# Define iteration variables as JAX arraysrc3NK|]}t|dr t|n|ywrrrs r)rz.PallasKernel.codegen_kernel..s'7KL'!Y*?AQ F7rrrc32K|]}t|ywr1rirs r)rz.PallasKernel.codegen_kernel..6s-Qc!f-Qz = jnp.arange(z ).reshape(r= _jit_wrapperrr rc32K|]}t|ywr1ri)rr?s r)rz.PallasKernel.codegen_kernel..Ms,LSV,Lrz,)z()zB@functools.partial(jax.jit, static_argnums=(0, 1), donate_argnums=z(out_shapes, out_dtypes, zout_specs = tuple(z& jax.ShapeDtypeStruct(shape, dtype)z3 for shape, dtype in zip(out_shapes, out_dtypes)z<# Calculate block_size as next power of 2 for Triton backendz0# Find maximum flattened size across all tensorsz max_size = 0zmax_size = max(max_size, z.size)zfor shape in out_shapes:zC tensor_size = shape[0] if len(shape) == 1 else math.prod(shape)z) max_size = max(max_size, tensor_size)z&block_size = next_power_of_2(max_size)Fc30K|]\}}|d|yw)z: Nr7)rr@os r)rz.PallasKernel.codegen_kernel..zs)P&1aQCr!+)Pszfunctools.partial(z _kernel, block_size=block_size),z_kernel,zreturn pl.pallas_call(z z out_shape=out_specs,z interpret=,z grid=(1,),z input_output_aliases={ z },z input_output_aliases={},z)(_mainz, stream=None):z/# Enable JAX x64 mode for float64/int64 supportz)jax.config.update('jax_enable_x64', True)z*# Convert Torch -> JAX for donated outputsz_jax = jax.device_put(z-.cpu().numpy(), device=jax.devices('tpu')[0])z_jax = jax.dlpack.from_dlpack(z .detach())z+# Convert Torch -> JAX for in-place tensorsz!# Convert Torch -> JAX for inputsrz.detach().contiguous())z-# Prepare output metadata from PyTorch tensorzout_shapes = (ztuple(rPzout_dtypes = (ztorch_dtype_to_jax_runtime(r_jaxc3(K|] }| ywr1r7)rr arg_name_maps r)rz.PallasKernel.codegen_kernel..s++/L&+szres = z(out_shapes, out_dtypes)z9result_values = res if isinstance(res, tuple) else (res,)z'res_cpu = jax.device_get(result_values[z])z".copy_(torch.from_dlpack(res_cpu))z'.copy_(torch.from_dlpack(result_values[z])))3rr.python_argdefsrr RuntimeErrorrrqrrrrrrr _inductorr _debug_cpu_to_tpu_pallas!pallas_take_first_jax_device_onlyr rsplicegetrrr4aliasable_out_ptrsrNrUindentrrr2rrrrrr5rrrr r!rr_linesrrrgetvalue)Br(rcodearg_defsrr| kernel_paramsrgpure_out_params output_paramsouterinneroutput_buffer_lookup kernel_nameinterpret_is_cpuis_tpuinterpret_literalimportsaliasable_flagsparamrr alias_params pointer_tailkernel_input_paramsfull_kernel_paramsflagnon_alias_out_setidxcopy_output_indiceskernel_signature buf_to_paramrr param_namematching_paramfirst_output_shapefirst_output_numelfirst_out_paramfirst_out_buf_namerrrvar_symrmrAr5 length_str length_val shape_strliner store_linejit_wrapper_namedonate_indicesdonate_literal alias_pairsout_idx alias_name input_idxalias_map_literal kernel_arg main_nameptralias_args_strrsB @r)codegen_kernelzPallasKernel.codegen_kernels !II446!Q)12A2 2&3Oq||I7N1OO$  5N(OA  ST T!% 8 8 > > @ u%% 5L  -o 77>>@EEN''@@ ??))KK"X"#"k'7FG  /3.A.A*r K (W "&&b%  , G4 (+-$ NE.2259K't38R8R9M+;&:%M OE "  N+: !&_U=SugV   %  5M(NA  +\9)M9&$3$9$9$; HjdD4T H "+=!9 CTEV=VC  #2;-x 2D(E'F G$($7$7 R A   '( [[]Z /""t'8'8! $(II$;$;$A$A$CULE53=eS3I%uL'U$(II$<$<$B$B$DULE53=eS3I%uL'U+11B1B1H1H1J*K&Hh!-!1!1(H!EJ%)N*"%?h#a&.@-.N! " &XJ'?@(8N;K5'QR'j(DXJeT,$$T-@-@KL%)"%)" &3A&6O)=)A)A/)R&) !"#''"4"45G"HC#&<<>D167PT72.23.%78 2a 7 28 '+&;&;&A&A&CQNGU"7|H"\\F!%F!3J*4;FI4NS[TX ") + 23a7&*<<$(II-Q>P-Q$Q 'jzl*YKWXY(>*Q'OP5Q: ++ *s4y) * (,'='= /#00NN:. /qZ /x*],7"#67 /IC $)F%%cAg. /  499,L^,L#LLtSN!N  >a ! #$$=diiH[>\=]]_ ` [[]8 NN/ 0 NNC D NNP Q NN3 ""RQR~.0NENN%>ugV#LMN9:YJKGH13K!*=!9 = ??9-&**47(,vV_ $7$=$=j$I #**Iw+?@ 3 9 9$ ?I&& 7';< =!% )PK)P P  &&%[M1QR#}H-  NN3 4 NN6J. / NN5 6 NN^,=+>a@ A NN+ , NN//@.AF3  NN4 "dii0C&D%EQGH NN3 q8 t#m5)  9+Qtyy);<=_ M [[]V  NNL M NNF GKL". J )l*@ Lyz)l*H T^_  NNH I# >>,/"e#9#>kl"e#A#jQ  NN> ?# >>(+"e#9#>kl"e#A#F]^  NNJ K NN )) NvdV73NOP  NN ))%2 6dV7C ,.L* ? .8\+> Z( ?# 1'*e4L S! 1#!%+3F+"-..GGWWXY(8'99QRS"O/ C(-DEcU"M$/Q'RS#f$KC5PST YV p}}q 3O  l  I ~ )! !&z2*%) *}Z /Z /Z8 8 TO eV p}}s ooo)o o :o; ooo*o o o 7 o&o&C7ppB p,Ao,;A po<B/ppDp!8D!)Q.67aff77&8TALL.C>-D) rrrrrrrrrNone)rrrr)rrrr)rrrztuple[str, bool])rr)rrrr)rrrrrrr1) rrrrrVrrWrrr) rrrrrtrrV+Union[CSEVariable, tuple[CSEVariable, ...]]rr)rrrr)rr2rr)rrrzOptional[IRNode]rr)r3r4r5r6r; overridesrr__annotations__r*rrrrrrrr rrrr\rrrrr __classcell__)rs@r)rrs &I).E &. :88&089=8FJ8 8>>T1l  . 5$=~+ Bm`MQ4949 *493>49FI49 49lDDD& D ; D 5 DL&& kZ ''r+rc:eZdZeZeddZ ddZy)PallasSchedulingc6ttjgSr1)rrREDUCE_TO_SINGLE_ELEMENT)clsrs r)get_backend_featuresz%PallasScheduling.get_backend_featuress>BBCDDr+ctjj}||jvr|j|Stj j r$t|tj j nd}tj|jdjdd}|dk(rd|}nd|d|}||j|<|jd|}t}|jd|d |j|d |jd t!||\} } | d | } |j#||j%| |S)Nr>zutf-8fusedpallas_rrzasync_compile.pallas(z, r'''Trz''') )rrr src_to_kernelr tritondescriptive_namesrhashlibsha256encode hexdigestr8rrUrr define_kernelr) r(src_code node_schedulerr fused_name kernel_hashrcompile_wrapperoriginsdetailed_originsmetadata_comments r)r zPallasScheduling.define_kernelsZ ''&& w,, ,((2 2}}.. "-1P1P Q  nnX__W%=>HHJ2AN  #K=1K#JK*5h'##O[A(*!!$9+"OPxt4!!&)$7 w$O!!%Yb)9(:;k?+C+C+EGWXr+N)rz torch.devicerzOrderedSet[BackendFeature])r rrzSequence[BaseSchedulerNode]rrrr)r3r4r5r kernel_type classmethodrr r7r+r)rrsFKEE ""3" " "r+r)7 __future__rrrtypingrrrrrrtorch.utils._ordered_setrtorch.utils._pallasr r>r runtime.runtime_utilsr rrr virtualizedrblock_analysisrcommonrrrrsimdrrrcollections.abcrrirr ops_handlerr schedulerr MAIN_SUFFIX_logginggetArtifactLoggerr3r&r"rr9r;rrr7r+r)r's" 66 /.6>/LL332+- ..228]K>>2U,U_ ,K_ ,D^':^'B!+~+r+