+ j^RIHt^RIt^RIt^RIt^RIt^RIt^RIHtH t H t ^RI t ^RI t ^RI Ht^RIHtHtHt^RIHt^RIHt^RIHtHt^R IHt^R IHt^R I H!t!H"t"H#t#H$t$H%t%^R I&H't'H(t(H)t)!R R]%4t*]*!4PVt,] 'd^RI-H.t.H/t/^RI0H1t1^RI2H3t3^RI4H5t5Rt6] PnPq]9R4t:!RR4t;!RR]<4t=!RR]$4t>]P~!RR44t@]P~!RR44tA]P~!RR 44tB]P~!R!R"44tC!R#R$](4tD!R%R&])4tER#)') annotationsN)Anycast TYPE_CHECKING) OrderedSet)MaxMinModularIndexing)configtorch_dtype_to_jax)get_fused_kernel_nameget_kernel_metadata)V)BlockPatternMatcher)BackendFeature CSEVariableIndentedBuffer OpOverrides PythonPrinter)IterationRangesEntry SIMDKernelSIMDSchedulingc>]tRt^#tRtRRltRRltRRltRtR #) PallasPrinterzG Custom sympy printer for Pallas that handles JAX-specific constructs. c V^8dQhRRRR/#expr sympy.Exprreturnstr)formats"_/home/luca/screen-checking/.venv/lib/python3.14/site-packages/torch/_inductor/codegen/pallas.py __annotate__PallasPrinter.__annotate__(s++++c VPVP^,4pVPVP^,4pVPVP^,4pRV RV RV R2#)z!Convert sympy Where to jnp.where. jnp.where(, ))doprintargs)selfrcpqs&& r% _print_WherePallasPrinter._print_Where(s_ LL1 & LL1 & LL1 &A3b2aS**r(c V^8dQhRRRR/#rr#)r$s"r%r&r'/zcr(c VPUu.uFq PV4NK ppV^,pVR,F pRV RV R2pK V#uupi)z7Convert sympy Min to jnp.minimum for JAX compatibility.NN jnp.minimum(r+r,r.r-r/rargr.results&& r% _print_MinPallasPrinter._print_Min/X-1YY7Yc S!Y7a88C#F82cU!4F 8Ac V^8dQhRRRR/#rr#)r$s"r%r&r'7r6r(c VPUu.uFq PV4NK ppV^,pVR,F pRV RV R2pK V#uupi)z7Convert sympy Max to jnp.maximum for JAX compatibility.r8 jnp.maximum(r+r,r;r<s&& r% _print_MaxPallasPrinter._print_Max7rArBr#N) __name__ __module__ __qualname____firstlineno____doc__r3r?rF__static_attributes__r#r(r%rr#s+r(r)CallableSequence)IRNode) ReductionType)BaseSchedulerNodemain kernel_codec8]tRt^StRtRRRlltRR/RltRtR#) PallasKernelWrapperz6Wrapper to provide .run() interface for Pallas kernelsNc V^8dQhRRRR/#)r kernel_fnzCallable[..., Any] kernel_path str | Noner#)r$s"r%r& PallasKernelWrapper.__annotate__VsDD"4D:Dr(c JWnW n\PRV4R#)zPallas kernel path: %sN)rXrYkernel_code_loginfo)r/rXrYs&&&r%__init__PallasKernelWrapper.__init__Vs"&5{Cr(streamc *VP!VRV/VB#)z Execute the Pallas kernel. Args: *args: Arguments to pass to the kernel function stream: CUDA stream to pass to the kernel function **kwargs: Additional keyword arguments for the kernel Returns: Result of the kernel execution ra)rX)r/rar.kwargss&$*,r%runPallasKernelWrapper.run[s~~t=F=f==r()rXrYN)rHrIrJrKrLr_rdrMr#r(r%rVrVSs@D > > >r(rVc]tRt^jtRtRtR#) UnsupportedzJException raised when an operation is not supported by the Pallas backend.r#N)rHrIrJrKrLrMr#r(r%rhrhjsTr(rhc ]tRt^ntRt]RRl4t]RRl4t]RRl4t]RR l4t ]R R l4t ]R R l4t ]RRl4t ]RRl4t ]RRl4t]RRl4t]RRl4t]RRl4t]RRl4t]RRl4t]RRl4t]R R!l4t]R"R#l4t]R$R%l4t]R&R'l4t]R(R)l4t]R*R+l4t]R,R-l4t]R.R/l4t]R0R1l4t]R2R3l4t]R4R5l4t]R6R7l4t ]R8R9l4t!]R:R;l4t"]R<R=l4t#]R>R?l4t$]RRARBll4t%]RCRDl4t&]RERFl4t']RGRHl4t(]RIRJl4t)]RKRLl4t*]RMRNl4t+]RORPl4t,]RQRRl4t-]RSRTl4t.]RURVl4t/]RWRXl4t0]RYRZl4t1]R[R\l4t2]R]R^l4t3]R_R`l4t4]RaRbl4t5]RcRdl4t6]ReRfl4t7]RgRhl4t8]RiRjl4t9]RkRll4t:]RmRnl4t;]RoRpl4t<]RqRrl4t=]RsRtl4t>]RuRvl4t?]RwRxl4t@]RyRzl4tA]R{R|l4tB]R}R~l4tC]CtD]RRl4tE]RRl4tF]RRl4tG]RRl4tH]RRl4tI]RRl4tJ]RRl4tK]RRl4tL]RRl4tM]RRl4tN]RRl4tO]RRl4tP]NtQ]RRl4tR]OtS]RRl4tT]RRl4tU]RRl4tV]UtW]VtX]RRl4tY]RRl4tZ]RRl4t[]RRl4t\]RRl4t]]RRl4t^]RRl4t_]RRl4t`]RRl4ta]RRl4tb]RRl4tc]RRl4td]RRl4te]RRl4tf]RRl4tg]RRl4th]RRl4ti]RRl4tj]RRl4tk]RRl4tl]RRl4tm]RRl4tn]RRl4to]RRl4tp]RRl4tq]RRl4tr]RRl4ts]RRl4tt]RRl4tu]RRl4tv]RRl4tw]RRl4tx]RRl4ty]RRl4tz]RRl4t{Rt|R@#)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 V^8dQhRRRR/#rxr"r!r#)r$s"r%r&"PallasKernelOverrides.__annotate__xssr(c RV R2#)zjnp.sin(r,r#rms&r%sinPallasKernelOverrides.sinv!Ar(c V^8dQhRRRR/#rlr#)r$s"r%r&rn}ror(c RV R2#)zjnp.cos(r,r#rqs&r%cosPallasKernelOverrides.cos{rtr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnror(c RV R2#)zjnp.tan(r,r#rqs&r%tanPallasKernelOverrides.tanrtr(c V^8dQhRRRR/#rlr#)r$s"r%r&rn    r(c RV R2#)z jnp.sinh(r,r#rqs&r%sinhPallasKernelOverrides.sinh1#Qr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnr~r(c RV R2#)z jnp.cosh(r,r#rqs&r%coshPallasKernelOverrides.coshrr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnr~r(c RV R2#)z jnp.tanh(r,r#rqs&r%tanhPallasKernelOverrides.tanhrr(c V^8dQhRRRR/#rlr#)r$s"r%r&rn""""r(c RV R2#)z jnp.arcsin(r,r#rqs&r%asinPallasKernelOverrides.asinQCq!!r(c V^8dQhRRRR/#rlr#)r$s"r%r&rnrr(c RV R2#)z jnp.arccos(r,r#rqs&r%acosPallasKernelOverrides.acosrr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnrr(c RV R2#)z jnp.arctan(r,r#rqs&r%atanPallasKernelOverrides.atanrr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnror(c RV R2#)zjnp.exp(r,r#rqs&r%expPallasKernelOverrides.exprtr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnr~r(c RV R2#)z jnp.exp2(r,r#rqs&r%exp2PallasKernelOverrides.exp21#Qr(c V^8dQhRRRR/#rlr#)r$s"r%r&rn!!!!r(c RV R2#)z jnp.expm1(r,r#rqs&r%expm1PallasKernelOverrides.expm1A3a  r(c V^8dQhRRRR/#rlr#)r$s"r%r&rnror(c RV R2#)zjnp.log(r,r#rqs&r%logPallasKernelOverrides.logrtr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnrr(c RV R2#)z jnp.log10(r,r#rqs&r%log10PallasKernelOverrides.log10rr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnr~r(c RV R2#)z jnp.log2(r,r#rqs&r%log2PallasKernelOverrides.log2rr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnrr(c RV R2#)z jnp.log1p(r,r#rqs&r%log1pPallasKernelOverrides.log1prr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnr~r(c RV R2#)z jnp.sqrt(r,r#rqs&r%sqrtPallasKernelOverrides.sqrtrr(c V^8dQhRRRR/#rlr#)r$s"r%r&rns%%%%r(c RV R2#)zjax.lax.rsqrt(r,r#rqs&r%rsqrtPallasKernelOverrides.rsqrts s!$$r(c V^8dQhRRRR/#rlr#)r$s"r%r&rnror(c RV R2#)zjnp.abs(r,r#rqs&r%absPallasKernelOverrides.absrtr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnsssr(c RV R2#)z(-r,r#rqs&r%negPallasKernelOverrides.negsA3ayr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnrr(c RV R2#)z jnp.floor(r,r#rqs&r%floorPallasKernelOverrides.floorA3a  r(c V^8dQhRRRR/#rlr#)r$s"r%r&rnr~r(c RV R2#)z jnp.ceil(r,r#rqs&r%ceilPallasKernelOverrides.ceilrr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnrr(c RV R2#)z jnp.trunc(r,r#rqs&r%truncPallasKernelOverrides.truncrr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnrr(c RV R2#)z jnp.round(r,r#rqs&r%roundPallasKernelOverrides.roundrr(c V^8dQhRRRR/#rlr#)r$s"r%r&rns&&3&3&r(c RV R2#)zjax.nn.sigmoid(r,r#rqs&r%sigmoidPallasKernelOverrides.sigmoid 1%%r(c V^8dQhRRRR/#rlr#)r$s"r%r&rns&&&&r(c RV R2#)rEz, 0)r#rqs&r%reluPallasKernelOverrides.relusaS%%r(c$V^8dQhRRRRRR/#rar"br!r#)r$s"r%r&rns!&&s&s&s&r(c RV RV R2#)z jnp.power(r+r,r#rrs&&r%powPallasKernelOverrides.powsA3b1%%r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rn!((3(3(3(r(c RV RV R2#)rEr+r,r#rs&&r%maximumPallasKernelOverrides.maximumaS1#Q''r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rnrr(c RV RV R2#)r:r+r,r#rs&&r%minimumPallasKernelOverrides.minimumrr(c(V^8dQhRRRRRRRR/#)rcondr"rrr!r#)r$s"r%r&rns(..C.C.C.C.r(c RV RV RV R2#)r*r+r,r#)rrrs&&&r%wherePallasKernelOverrides.wheresD6A3b1--r(c(V^8dQhRRRRRRRR/#)rmaskr"bodyzCallable[[], str]otherfloatr!r#)r$s"r%r&rns);;S; 1;%;C;r(c V!4p\V\4'dT\P!V4'dRpM?\P!V4'd V^8dRMRpM\ V4pM \ V4pRV RV RV R2#)zr Computes body, but only uses the result where mask is true. Where mask is false, uses the 'other' value instead. jnp.nanjnp.inf-jnp.infr*r+r,) isinstancermathisnanisinfrepr)rrrr> other_strs&&& r%maskedPallasKernelOverrides.maskedsu  eU # #zz%  % E"").I  K U ID6F82i[::r(Nc ,V^8dQhRRRRRRRRR R/#) rrmr"dtype torch.dtype src_dtypeztorch.dtype | Noneuse_compute_typesboolr!r#)r$s"r%r&rns: 7 7  7 7& 7 7 7r(c V\P8Xd>\PP 4P R8Xd\P p\V4pRV RV R2#)tpuz jnp.asarray( ).astype(r,)torchint64rgraphget_current_device_or_throwtypeint32r )rmr r r jax_dtypes&&&& r%to_dtypePallasKernelOverrides.to_dtypesS EKK AGG$G$G$I$N$NRW$WKKE&u- aS )A66r(c(V^8dQhRRRRRRRR/#)rrmr"r r r r!r#)r$s"r%r&rn&s/ffCf f fPSfr(c F\V4p\V4pRV RV RV R2#)z=Bitcast a value from one dtype to another with the same size.z)jax.lax.bitcast_convert_type(jnp.asarray(rz), r,r )rmr r r jax_src_dtypes&&& r%to_dtype_bitcast&PallasKernelOverrides.to_dtype_bitcast%s7'u- *95 :1#Y}oUXYbXccdeer(c$V^8dQhRRRRRR/#rrr r r r!r"r#)r$s"r%r&rn.s!:::K:C:r(c ^RIHp\PPP \PP V44\PPV4p\PPV4p\PPV4p\PPP\PPWR!V4R7p\PWa4#)z>Convert a sympy expression to a JAX array indexing expression.)get_bounds_index_expr)bounds)utilsr$rkernelused_iter_varsupdate_get_used_iter_varsprepare_indexingrename_indexingkexprcsegeneratecomputerjr)rr r$preparedrenamedidx_strvars&& r% index_expr PallasKernelOverrides.index_expr-s 2 &&qxx'C'CD'IJ88,,T2((**84((..)hhll## HH  g.CD.I$ %--c99r(c$V^8dQhRRRRRR/#r"r#)r$s"r%r&rn?s!===K=C=r(c ,\PW4#rf)rjr5)rr s&&r% value_expr PallasKernelOverrides.value_expr>s%//< r,r#rs&&r%gtPallasKernelOverrides.gt}rjr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnrr(c RV R2#)z jnp.isnan(r,r#rqs&r%rPallasKernelOverrides.isnanrr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnrr(c RV R2#)z jnp.isinf(r,r#rqs&r%rPallasKernelOverrides.isinfrr(c V^8dQhRRRR/#rlr#)r$s"r%r&rns$$C$C$r(c RV R2#)z jnp.isfinite(r,r#rqs&r%isfinitePallasKernelOverrides.isfinitesqc##r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rnr\r(c RV RV R2#)rXz >= r,r#rs&&r%gePallasKernelOverrides.ger`r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rn!,,s,s,s,r(c RV RV R2#)zjnp.logical_and(r+r,r#rs&&r% logical_and!PallasKernelOverrides.logical_and"!Bqc++r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rn!++c+c+c+r(c RV RV R2#)zjnp.logical_or(r+r,r#rs&&r% logical_or PallasKernelOverrides.logical_ors!2aS**r(c V^8dQhRRRR/#rlr#)r$s"r%r&rn''s's'r(c RV R2#)zjnp.logical_not(r,r#rqs&r% logical_not!PallasKernelOverrides.logical_not!!A&&r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rnrr(c RV RV R2#)zjnp.logical_xor(r+r,r#rs&&r% logical_xor!PallasKernelOverrides.logical_xorrr(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rns!(((((r(c RV RV R2#)z jnp.arctan2(r+r,r#rs&&r%atan2PallasKernelOverrides.atan2rr(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rn!&&&&&r(c RV RV R2#)z jnp.hypot(r+r,r#rs&&r%hypotPallasKernelOverrides.hypotA3b1%%r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rns!%%%%%r(c RV RV R2#)z jnp.fmod(r+r,r#rs&&r%fmodPallasKernelOverrides.fmods1#Rs!$$r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rn!**S*S*S*r(c RV RV R2#)zjnp.remainder(r+r,r#rs&&r% remainderPallasKernelOverrides.remainderss"QCq))r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rns&eeCeCeCer(c  &RV RV RV RV RV R2 #)z (jnp.sign(z ) * jnp.sign(z ) * (jnp.abs(z ) // jnp.abs(z ))).astype(.dtype)r#rs&&r%truncdivPallasKernelOverrides.truncdivs. A3mA3mA3mA3kZ[Y\\cddr(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rns!CCCr(c RV RV R2#)rXz // r,r#rs&&r%floordivPallasKernelOverrides.floordivr`r(c(V^8dQhRRRRRRRR/#)rrmr"min_valmax_valr!r#)r$s"r%r&rns(666s6S6S6r(c RV RV RV R2#)z jnp.clip(r+r,r#)rmrrs&&&r%clampPallasKernelOverrides.clamps1#Ry7)155r(c V^8dQhRRRR/#rlr#)r$s"r%r&rns@@@@r(c RV RV R2#)zjnp.where(jnp.isnan(z), 0.0, jnp.sign())r#rqs&r%signPallasKernelOverrides.signs&aS(9!B??r(c V^8dQhRRRR/#rlr#)r$s"r%r&rns##3#3#r(c RV R2#)z jnp.signbit(r,r#rqs&r%signbitPallasKernelOverrides.signbitsaS""r(c V^8dQhRRRR/#rlr#)r$s"r%r&rns--s-s-r(c RV R2#)zjax.scipy.special.erf(r,r#rqs&r%erfPallasKernelOverrides.erfs(s!,,r(c V^8dQhRRRR/#rlr#)r$s"r%r&rns....r(c RV R2#)zjax.scipy.special.erfc(r,r#rqs&r%erfcPallasKernelOverrides.erfcs(1--r(c V^8dQhRRRR/#rlr#)r$s"r%r&rns00#0#0r(c RV R2#)zjax.scipy.special.erfinv(r,r#rqs&r%erfinvPallasKernelOverrides.erfinvs+1#Q//r(c V^8dQhRRRR/#rlr#)r$s"r%r&rns11#1#1r(c RV R2#)zjax.scipy.special.gammaln(r,r#rqs&r%lgammaPallasKernelOverrides.lgammas,A3a00r(c V^8dQhRRRR/#rlr#)r$s"r%r&rns113131r(c RV R2#)zjax.scipy.special.digamma(r,r#rqs&r%digammaPallasKernelOverrides.digammas+A3a00r(c V^8dQhRRRR/#rlr#)r$s"r%r&rn   S  S  r(c RV RV RV R2#)r*z>.astype(jnp.float64) == 0.0, 1.0, jax.scipy.special.bessel_jn(z&.astype(jnp.float64), v=0)[0]).astype(rr#rqs&r% bessel_j0PallasKernelOverrides.bessel_j0+++,#.c " r(c V^8dQhRRRR/#rlr#)r$s"r%r&rnrr(c RV RV RV R2#)r*z>.astype(jnp.float64) == 0.0, 0.0, jax.scipy.special.bessel_jn(z&.astype(jnp.float64), v=1)[1]).astype(rr#rqs&r% bessel_j1PallasKernelOverrides.bessel_j1rr(c V^8dQhRRRR/#rlr#)r$s"r%r&rn BBcBcBr(c RV RV R2#)jax.lax.bessel_i0e() * jnp.exp(jnp.abs(rr#rqs&r%modified_bessel_i0(PallasKernelOverrides.modified_bessel_i0%QC';A3bAAr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnrr(c RV RV R2#)jax.lax.bessel_i1e(rrr#rqs&r%modified_bessel_i1(PallasKernelOverrides.modified_bessel_i1rr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnsAAsAsAr(c RV RV RV R2#)r*z == 0.0, 1.0, jnp.sin(z) / r,r#rqs&r%spherical_bessel_j0)PallasKernelOverrides.spherical_bessel_j0s A34QCtA3a@@r(c V^8dQhRRRR/#rlr#)r$s"r%r&rn**s*s*r(c RV R2#)rr,r#rqs&r%i0ePallasKernelOverrides.i0e%QCq))r(c V^8dQhRRRR/#rlr#)r$s"r%r&rn$rr(c RV R2#)rr,r#rqs&r%i1ePallasKernelOverrides.i1e#rr(c$V^8dQhRRRRRR/#rrmr"yr!r#)r$s"r%r&rn)s!77C7C7C7r(c RV RV R2#)zjax.scipy.special.gammainc(r+r,r#rmrs&&r%gammaincPallasKernelOverrides.gammainc(s-QCr!A66r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rn/s!88S8S8S8r(c RV RV R2#)zjax.scipy.special.gammaincc(r+r,r#rs&&r% gammainccPallasKernelOverrides.gammaincc.s.aS1#Q77r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rn8s&JJSJSJSJr(c RV RV R2#)zjax.scipy.special.polygamma(z.astype(jnp.int32), r,r#rs&&r% polygammaPallasKernelOverrides.polygamma7s.aS0DQCqIIr(c V^8dQhRRRR/#rlr#)r$s"r%r&rn>s////r(c RV R2#)zjax.scipy.special.ndtri(r,r#rqs&r%ndtriPallasKernelOverrides.ndtri=s*!A..r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rnCs!33333r(c RV RV R2#)zjax.scipy.special.zeta(r+r,r#rs&&r%zetaPallasKernelOverrides.zetaBs)2aS22r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rnHs!44444r(c RV RV R2#)zjax.scipy.special.xlogy(r+r,r#rs&&r%xlogyPallasKernelOverrides.xlogyGs*!Bqc33r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rnMs!66363636r(c RV RV R2#)zjax.scipy.special.xlog1py(r+r,r#rs&&r%xlog1pyPallasKernelOverrides.xlog1pyLs,A3b155r(c$V^8dQhRRRRRR/#rrmr"nr!r#)r$s"r%r&rnSs!   #  #  #  r(c >RV RV RV RV RV RV RV RV RV R 2#) 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#rmrs&&r%chebyshev_polynomial_t,PallasKernelOverrides.chebyshev_polynomial_tQs[!$c04s5aS9M!,H:  W r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rnbs!  # # # r(c RP.RNV NRNV NRNV NRNV NRNV NRNV NRNV NRNV NR NV NR NV NR NV NR NV NR NV NRNV NRNV NRNV NR NV NRN4#)rz) < 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)))))joinrs&&r%chebyshev_polynomial_u,PallasKernelOverrides.chebyshev_polynomial_u`s ; ; ; ;$ ;s ;6 ;78c ;:( ;()s ;+ ; ; ; ;# ;$%# ;& ;  ; < ; =>3 ; ?$ ; %&3 ; ' ; ;/ ;01c ;27 ;89c ;: ; ;+ ;,-# ;.O ;PQc ;R$ ;%&3 ;': ; r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rnx!   #  #  #  r(c RP.RNV NRNV NRNV NRNV NRNV NRNV NRNV NRNV NRNV NR NV NR NV NR NV NR NV NR NV NRNV NRNV NR NV NRNV NRNV NRNV NRNV NRNV NRNV NRN4#)r#r* == 0, jnp.ones_like( ), jnp.where( == 1, 2* - 1, jnp.where( == 2, 4*z**2 - 2* == 3, 8*z**3 - 4***2 - 4* + 1, jnp.where( == 4, 16*z**4 - 8* **3 - 12*z**2 + 4* == 5, 32*z **5 - 16* **4 - 32*z **3 + 12***2 + 6*z - 1, jnp.zeros_like()))))))r&rs&&r%chebyshev_polynomial_v,PallasKernelOverrides.chebyshev_polynomial_vv  ) )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(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rnr+r(c RP.RNV NRNV NRNV NRNV NRNV NRNV NRNV NRNV NR NV NR NV NR NV NRNV NR NV NR NV NRNV NR NV NRNV NRNV NRNV NRNV NRNV NRNV NRNV NRN4#)r#r*r-r.r/r4r1z**2 + 2*r0r2z**3 + 4*r3r5z**4 + 8*r6r7z **5 + 16*r8r9z + 1, jnp.zeros_like(r:r&rs&&r%chebyshev_polynomial_w,PallasKernelOverrides.chebyshev_polynomial_wr=r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rn&QQ#Q#Q#Qr(c 6\PRV R2V4#z(2 * z - 1))rjrrs&&r%shifted_chebyshev_polynomial_t4PallasKernelOverrides.shifted_chebyshev_polynomial_t %;;eA3e3 )? )  )$  )%&3  )'0  )12s  )3<  )=>3  )?H  )IJs  )K  )   ) %  ) &'C  ) (1  ) 23  ) 4>  ) ?@S  ) AK  ) LM#  ) NX  ) YZWZ  ) [  )  S  ) !(  ) r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rnr\r(c RP.RNV NRNV NRNV NRNV NRNV NRNV NRNV NRNV NR NV NR NV NR NV NR NV NR NV NRNV NRNV NRNV NRNV NRN4#)r#r*r-r.r$rWz == 2, (3 * z**2 - 1) / 2, jnp.where(z == 3, (5 * r^z) / 2, jnp.where(z == 4, (35 * z **4 - 30 * z**2 + 3) / 8, jnp.where(z == 5, (63 * z **5 - 70 * r_z) / 8, jnp.zeros_like(r:r&rs&&r%legendre_polynomial_p+PallasKernelOverrides.legendre_polynomial_ps~  ) )j ) )0 ) )4 ) )" )#$# )& ) )' )()s )+ ) )( ))*s )+5 )67C )8 )  ))  )*+  ),7  )89c  ):  )   ) )  ) *+  ) ,7  ) 89c  ) :E  ) FGC  ) H  )  S  ) !(  ) r(c V^8dQhRRRR/#rlr#)r$s"r%r&rns&&c&c&r(c RV R2#)zjnp.reciprocal(r,r#rqs&r% reciprocal PallasKernelOverrides.reciprocalrr(c V^8dQhRRRR/#rlr#)r$s"r%r&rns""#"#"r(c RV R2#)z jnp.square(r,r#rqs&r%squarePallasKernelOverrides.squaresQCq!!r(c(V^8dQhRRRRRRRR/#)rrr"rr0r!r#)r$s"r%r&rns(**s*s*s*s*r(c RV RV RV R2#)zFused multiply-add: a * b + c JAX doesn't have jnp.fma, so we use the unfused version. The compiler may still fuse this on supported hardware. z((r%z) + (rr#)rrr0s&&&r%fmaPallasKernelOverrides.fmasA3eA3eA3b))r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rns!))C)C)C)r(c RV RV R2#)z jnp.copysign(r+r,r#rs&&r%copysignPallasKernelOverrides.copysignsqcA3a((r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rnrr(c RV RV R2#)zjnp.nextafter(r+r,r#rs&&r% nextafterPallasKernelOverrides.nextafters s"QCq))r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rn rr(c RV RV R2#)z jnp.ldexp(r+r,r#rs&&r%ldexpPallasKernelOverrides.ldexp rr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnrr(c RV R2#)z jnp.frexp(r,r#rqs&r%frexpPallasKernelOverrides.frexprr(c V^8dQhRRRR/#rlr#)r$s"r%r&rnr~r(c RV R2#)z jnp.modf(r,r#rqs&r%modfPallasKernelOverrides.modfrr(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rnrr(c RV RV R2#)zjnp.bitwise_and(r+r,r#rs&&r% bitwise_and!PallasKernelOverrides.bitwise_and!!Bqc++r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rnrr(c RV RV R2#)zjnp.bitwise_or(r+r,r#rs&&r% bitwise_or PallasKernelOverrides.bitwise_or 2aS**r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rn!rr(c RV RV R2#)zjnp.bitwise_xor(r+r,r#rs&&r% bitwise_xor!PallasKernelOverrides.bitwise_xor rr(c V^8dQhRRRR/#rlr#)r$s"r%r&rn%rr(c RV R2#)zjnp.bitwise_not(r,r#rqs&r% bitwise_not!PallasKernelOverrides.bitwise_not$rr(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rn)rr(c RV RV R2#)zjnp.left_shift(r+r,r#rs&&r% left_shift PallasKernelOverrides.left_shift(rr(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rn-rr(c RV RV R2#)zjnp.right_shift(r+r,r#rs&&r% right_shift!PallasKernelOverrides.right_shift,rr(c$V^8dQhRRRRRR/#)rnamer"offsetr!r#)r$s"r%r&rn2s&CCCSCSCr(c \PPPRV4pR\PPP V4 RV R2#)z)Load the random seed value from a buffer.load_seed_offsetrXz[0] + r,)rr'r. seed_offsetinput)rrrs&& r% load_seedPallasKernelOverrides.load_seed1sH hhmm//0BFK 188==&&t,-VK=BBr(c$V^8dQhRRRRRR/#rseedr"rr!r#)r$s"r%r&rn:s!   3      r(c RV RV RV R2#)zGenerate 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#rrs&&r%randPallasKernelOverrides.rand9s,@@DvF"8#DVHH V r(c$V^8dQhRRRRRR/#rr#)r$s"r%r&rnJs!   C      r(c RV RV RV R2#)zGenerate 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(rrrr#rs&&r%randnPallasKernelOverrides.randnIs,@@DvF"8#DVHH V r(c ,V^8dQhRRRRRRRRRR/#)rrr"rlowhighr!r#)r$s"r%r&rnYs/   S s # # r(c  &RV RV RV RV RV R2 #)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(rrr#)rrrrs&&&&r% randint64PallasKernelOverrides.randint64Xs> @@DvE]^a]bbdeidjk"8#DVHH V r(r#)NT)}rHrIrJrKrL staticmethodrrrwr{rrrrrrrrrrrrrrrrrrrrrrrrrrrrrrr5r9r@rDrHrLrPrTrYr^rcrhrmrqrrr{rrrrrrrrrrrrcliprrrrrrrrrrrri0ri1rrrigammaigammacrr rrrrr(r;r@rFrKrOrSrYr`rdrhrlrprtrxr|rrrrrrrrrrrrrrMr#r(r%rjrjns       """"""  !!!!  !!  %%!!  !!!!&&&&&&((((..;;& 7 7ff:: ==  6 6      !!EE33 !!!!$$,,++'',,((&&%%**ee 66 D@@##--..001111        BB BB AA B** B**77 88FGJJ //334466      (        QQQQQQQQ  "            &&""**))**&&!!  ,,++,,''++,,CC          r(rjcX]tRtRt$RtR]R&R]R&R]R&R]R &R]R &R]R &R tR #)_IndirectAccessInfoicz=Describes a detected indirect (data-dependent) buffer access.r" table_paramtable_buf_nametuple table_shapeint indirect_dim indirect_var indices_paramr#NrHrIrJrKrL__annotations__rMr#r(r%rrcs)Gr(rc0]tRtRt$RtR]R&R]R&RtR#) _BufferIndexingiozHEncapsulates index string and flattening requirements for buffer access.r" index_strr needs_flattenr#Nrr#r(r%rrosRNr(rcD]tRtRt$RtR]R&R]R&R]R&R ]R &R tR #) _BroadcastedIterVariwzFEncapsulates information needed to codegen a broadcasted iteration varridx sympy.Symbolvar_symrentry int | None length_valr#Nrr#r(r%rrwsP H  r(rc]tRtRt$RtR]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&RtR#) _CodegenContextiz@Bundles local state shared across codegen_kernel helper methods.rcoder" kernel_nameris_tpuinterpret_is_cpuinterpret_literal list[str] kernel_paramspure_out_params output_paramssize_var_paramszdict[str, str]output_buffer_lookupdict[str, bool]aliasable_flags alias_params pointer_tailkernel_input_paramsfull_kernel_paramszOrderedSet[str]non_alias_out_set list[int]copy_output_indiceslist[tuple[int, int]] alias_pairsr#Nrr#r(r%rrspJ  L(($$""!!&&""&&r(rca]tRtRt$Rt]t]tR] R&V3Rlt RRlt RR lt R R lt R R ltRRltRRltRRltRRltRRltRRlt]RRl4tRRlt]RRl4tR R!ltR"R#ltR$R%lt]R&R'l4tR(R)lt]R*R+l4t]R,R-l4t]R.R/l4t R0R1lt!R2R3lt"R4R5lt#R6R7lt$R8R9lt%R:R;lt&R<R=lt'R>R?lt(R@RAlt)RBRClt*RDRElt+RFRGlt,]RHRIl4t-RJRKlt.RLRMlt/RNROlt0RPRQlt1RRRSlt2RTRUlt3RVRWlt4RXRYlt5RZR[lt6R\R]lt7RR^R_llt8R`Ralt9]:PvRbRcl4t]RRhRill4t?RRjRkllt@RlRmltARnRoltBRpRqltC]RrRsl4tDRtRultERRvRwlltF]RxRyl4tGRzR{ltHR|R}ltIR~RltJRRltKRRltLRRltMRRltN]RRl4tORRltPRRltQRRltRRRltSRRltTRRltURRltVRRltW]RRl4tXRRRlltYRtZV;t[#) PallasKernelizPallas 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. zCallable[[sympy.Expr], str]r-c <\SV`!V/VB\PP 4pVP R8HVnVP R8HVnVP Vn.Vn /Vn \4Vn /Vn /VnRVnRVn.VnVP$P'4FEpVP(P*F(pVP"P-VP.4K* KG \4Vn\4VnRVn/Vn\4VnRVn/Vn/VnVP R8HVnR#)cudarNF) superr_rrrris_gpuruse_emit_pipelinestore_with_outputload_index_exprsroutputs_need_readpermuted_input_bufferscollapsed_reshape_inputscollapsed_output_shape_cpu_max_grid_product_output_buffer_namesfeaturesscheduler_nodes read_writeswritesappendrr(tile_relative_iter_varshas_flatten_indexingstrided_input_buffersflatten_indexed_buffersindirect_access _cse_to_param_param_to_graph_name)r/r.rcdevicesnodedep __class__s&*, r%r_PallasKernel.__init__sB $)&)446kkV+ kkU* "&8:792<,BD#DF%>B#15"02!]]224E((//))00:059C BL$$)! MO"9C $;?-/46!kkU* r(c ,V^8dQhRRRRRRRRRR/#) rrr sizelowerrupperr!Noner#)r$s"r%r&PallasKernel.__annotate__s4888&089=8FJ8 8r(c R#)z)Check array bounds for indirect indexing.Nr#)r/rrrrs&&&&&r% check_boundsPallasKernel.check_boundssr(c V^8dQhRRRR/#rindexr r!r"r#)r$s"r%r&rs!>!>J!>3!>r(c VPV4p\V\P4'dR#VP'd \ V4#VP V4#)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+rsympySymbol is_Integerr"_convert_to_jax_slice)r/r"prepared_indexs&& r%_get_index_strPallasKernel._get_index_strsT"..u5 nell 3 3  & & &~& &--n= =r(c V^8dQhRRRR/#r!r#)r$s"r%r&rsbb:b#br(c \VP'gR#VPV4pVP\4'd<VPP VP V44VPV4#\PPPV4pVP V4pVPP V4\V4^8Xd \V4#\V4^8XEd*\\V44p\ P"!W4p\ P$!WC4pVeW, p\PPPV4pV^8dVPV4#V^8XdR#V^8wdVPV4#\'V4pV^8dVPV4#TPT4 R2#W, p\PPPV4pV^8Xd WC8XdR#R#\V4^8dORpVF:p\ P"!W4p\ P$!WC4pV^8wgK8RpM V'dR#R#R# \(\*3dTPT4u#i;i)a  Convert a sympy index expression to JAX slice notation. Handles common patterns like: - stride*var -> ::stride - stride*var + offset -> offset::stride For more complex patterns, falls back to explicit indexing. Uses BlockPatternMatcher for robust pattern matching. r$z::1TF) range_treesr,hasr r(r)r*r-rrsizevarssimplifylenr"nextiterrget_subexpr_involving_symbolmatch_affine_block_exprr TypeError ValueError) r/r" used_varsr4var_exprstrider offset_valall_unit_strides && r%r("PallasKernel._convert_to_jax_slices[$$U+ 99_ % %    & &t'?'?'F G::e$ $  ))%0,,U3  ""9- y>Q u:  ^q tI'C+GGSH)@@OF!)))226:A:::e,,Q; Q;::e,,-!$VJ!A~#zz%00& **V,-S11)))226:Q;8? ,+^a #O .KKEW,DDXSQ;&+O !  A":.-::e,,-s,"J"J+*J+c V^8dQhRRRR/#r!r#)r$s"r%r&rbsZCr(c VPpVP4pW#,pWB8wd\RV 24hVPP V4VP V4pV#)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_varsrhr(r)r-)r/r"rA iter_varsr9rs&& r%_generate_strided_index$PallasKernel._generate_strided_indexbsr)) '') !,  $KE7S  ""9-JJu% r(cV^8dQhRR/#)rr!rr#)r$s"r%r&rs88 8r(c H\VPP44#)z*Get the set of iteration variable symbols.)rrange_tree_nodeskeys)r/s&r%rBPallasKernel._get_iter_varss$//44677r(c V^8dQhRRRR/#)rr"r r!rr#)r$s"r%r&rs::: :r(c DVPVP4,#)z4Get iteration variables used in an index expression.)rArBr/r"s&&r%r* PallasKernel._get_used_iter_varss!!D$7$7$999r(c V^8dQhRRRR/#rr"r r!rr#)r$s"r%r&rs5555r(c 6\VPV44#)z7Check if index expression contains iteration variables.)rr*rMs&&r%_has_iteration_vars PallasKernel._has_iteration_varssD,,U344r(c V^8dQhRRRR/#)rr"r r!zlist[sympy.Symbol]r#)r$s"r%r&rsKK K7IKr(c VPUu.uF'p\V4PR4'gK%VNK) up#uupi)zDGet list of indirect variable symbols (tmp*) in an index expression.tmp)rAr" startswith)r/r"ss&& r%_get_indirect_varsPallasKernel._get_indirect_varss5 --J-aQ1B1B51I-JJJs"??c V^8dQhRRRR/#rPr#)r$s"r%r&rs77 7t7r(c <\VPV44^8#)z6Check if index expression contains indirect variables.)r2rYrMs&&r%_has_indirect_varsPallasKernel._has_indirect_varss4**512Q66r(c$V^8dQhRRRRRR/#)rr"r rr"r!z!list[tuple[int, int, int]] | Noner#)r$s"r%r&rs)III'*I *Ir(c VPV4'gVP\4'dR#W P9dR#VP V4pVfR#VwrEpVUu.uFq`P V4NK pp\ ;QJdRV4F 'gK RM RM !RV44'dR#\\\,V4p\V4p V ^8XdR#VPV4p VPV4p V 'gR#\V 4Uu.uFp^^.NK p p/p \PP P#V4pV EF#p\$P&!W4p\$P(!VV4pVfR#VP V4pVeV^8:dR#Rp\V 4FEpV V,^8XdKVV V,,^8XgK*VV V,,pV^8gKCTpM VfR#VV V,,pV^8gVV,V,^8wdR#VV V,^&VW&\PP P#VV, 4pEK& VP V4pVfR#V^8dR#\V 4F=pV V,^8gKVV V,,V V,^&VV V,,pK? V^8wdR#\*;QJdRV 4F 'dK RM RM !RV 44'dR#.p^p\V 4F_pV V,wppVV,pVV,pVV,V,pVV8dR#VVV, ,pVP-VVV34Ka VP/V4wppVV8wdR#V P14FzwppWP29dR#VP VP2V,P44pVfR#VV,wp p!p"VV,V ,V", p#VV#8wgKyR# V#uupiuupi)aDecompose 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. Nc3B"TFqRJ;'gV^8*xK R#5irfr#.0rXs& r% 9PallasKernel._decompose_strided_access..s#>,=qDy""AF",=sTFc30"TF wrV^8HxK R#5ir9Nr#)rcrX_s& r%rdres)&$!qAv&s)r]r/r r_get_buffer_info _safe_intanyrlistrr2_c_contiguous_stridesr*rangerrr0r1rr5r6allr  _compute_output_numel_from_indexitemsrHlength)$r/r"rr^rhbuf_sizerXbuf_shape_or_none buf_shapendim c_stridesr9r> var_to_dim remainingr4r:coeff coeff_intdimdper_dim_strider< decomposedoutput_numel_expectedr; raw_offsetrskipn_blocks output_numel var_rangestride_d _offset_dskip_deffective_sizes$&&& r%_decompose_strided_access&PallasKernel._decompose_strided_accesss$  " "5 ) )UYY-G-G // /$$T* <#Q18@A1^^A.A 3>,=>333>,=> > >#DI/@A 9~ 19..y9 ,,U3 49;"?;aAq6;"?.0 GG$$--e4 C*GG WH'??#NE}u-I INC4[Q<1$y|+q0%.)A,%>N%*!{&)C.8N!Ys^n%D%I+F3KN!JO((11)h2FGI9@^^I.   >tA|a)Yq\9q ! ')A,6  ? 3)&)333)&) ) ) 24 !tA!' FJ&(F'D |v-Hx !X_ 4 !   vvt4 5??F a 0 0#((*HC///t'<' 0``, slices off leading blocks so the remaining elements align with the output. N = .reshape(r+r,c32"TF wrV^8xK R#5irNr#)rcrhrs& r%rd9PallasKernel._codegen_strided_reshapes..Ks:'JAqtax'sTFrrr) _param_to_buf_namer ri enumeraterjr r" writeliner'rk)r/rrparambuf_namestridesr^rhrsnew_shape_partsr}r;_offsetrr| slice_partsrs&&& r%_codegen_strided_reshapes&PallasKernel._codegen_strided_reshapes+sE..u5H83M3M#M00:G((2D|#' AAq)+O/8/A++FUnnX[1;A:#**3s}+=>#**3v;7#**3s840BgSy?1K0LAN3:':333:':::-/K18-!Q;'..s3'..TAX$qz3O'..s3 29 NNeWCwa +8N7Oq#QR?r(c V^8dQhRRRR/#)rrrr!list[int] | Noner#)r$s"r%r&rVsC4Dr(c *\VRR4!4pVfR#\VRR4pVe\V4V8wdR#.pVFNp\V\\P 34'd \V4MRpVfR#VP V4KP V#)z5Extract actual output buffer strides from its layout. get_layoutcR#rfr#r#r(r%6PallasKernel._get_actual_out_strides..Xr(Nr;)getattrr2rrr%Integerr )out_bufrlayout stride_rawrrXvs&& r%_get_actual_out_strides$PallasKernel._get_actual_out_stridesUs, =? >VXt4  ZA!5A$Qemm(<==A4Ay NN1   r(c V^8dQhRRRR/#)rorderedrlr!z dict | Noner#)r$s"r%r&rfs  T k r(c /p.pVFTpVPV,pVPpWb9d.W&&VPV4W&,PV4KV .p\V4FpVP W&,4K /p^p VFBpVP VPV,P 4p V fR#WV&W,p KD V#)aCompute 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. N)rHprefixr reversedextendrjrr) r/r prefix_groups prefix_orderrnoder1inner_to_outercoeffsrzszs && r%_compute_store_coeffs"PallasKernel._compute_store_coeffsfs*, "$ A((+D A%#%  ##A&   # #A &  ",'A  ! !-"2 3(A 5 5a 8 ? ?@Bz1I KE   r(c$V^8dQhRRRRRR/#)rrr"r"r r!ztuple[int, ...] | Noner#)r$s"r%r&rs)hhh *h hr(c VPV4pV'gR#VwrErFVUu.uFqpPV4NK pp\V4^8gRV9dR#\\\ ,V4p V'gR#VP V4p VPP4UU u.uF#wr{Wz9gK V P'dK!VNK% p pp \V 4\V 48wEd\V 4p \V 4^8XEdsVPVPV ^,,P4\P!V 48XEd&VPV 4pVPEFp\P P#V4pVfK)VP%4Uu.uFqpPV4NK pp\&;QJdRV4F 'gK RM RM !RV44'g\V4V 8wdKVP)VV 4pVfR#VP+VV4pVfR#\ V4\ \-V 448XdR#\/V4u# R#V Uu.uF<pVP1\P P2P5V4V4NK> pp\6;QJdRV4F 'dK RM RM !RV44'gR#\\\ ,V4p\V 4p VPV 4pVP9V 4pVEewVPEFep\P P#V4pVfK)VP%4Uu.uFqpPV4NK pp\&;QJdRV4F 'gK RM RM !RV44'g\V4V 8wdKVP)VV 4pVeTP+V Uu.uF pVV,NK upV4pVP+\ V4V4pVefVeb^.V ,p\-V 4FpVV,VVV,&K \ V4\ \-V 448XdR#\/V4u#R# R#uupiuup piuupiuupiuupiuupi)aReturn 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. Nc3("TFqRJxK R#5irfr#rbs& r%rd:PallasKernel._get_full_load_permutation..s8i9iTFc3\"TF"p\V\4;'dV^8xK$ R#5irrrrcr0s& r%rdrs&DA:a%//!a%/,,c3("TFqRJxK R#5irfr#rbs& r%rdrs4)QDy)r)rirjr2rrlrr*rHrq is_reductionrrrprodrmrrr get_bufferget_sizerkr_map_coeffs_to_dimsrnr_get_index_coefficientr0r1ror)r/rr"r^rhrs is_contiguousrX in_shape_rawin_shape iter_usederr in_stridesout_namer out_shapeactualpermr coeffs_rawr store_coeffs rt_to_outrt_to_inks&&& r%_get_full_load_permutation'PallasKernel._get_full_load_permutations'($$T*+/(Q3;<8aq)8 < | q DL$8"49l;,,U3 --335 5~ &'nn A5  w<3x= ( H A7|q T^^%%gaj188&8$&%"77A $ 9 9Hgg00:G ./s&GA:a%//!a%/rFTc3("TFqRJxK R#5irfr#rbs& r%rdrEs4m9mr)rirjr2rrlrr*rHrqrrrrr_group_dims_to_rangesrnrrr0r1rrorgetr rrrrrkr)(r/rr"r^rhrsrrXrrrrrrr ranges_rawranges in_groupscollapsed_in_stridesr;i simplified in_coeffs_raw in_coeffsin_stride_to_dim var_to_in_dimrzr|rrr out_shape_rawr out_groupscollapsed_out_stridesjout_stride_to_dimvar_to_out_dimrrs(&&& r%_get_collapsed_load_permutation,PallasKernel._get_collapsed_load_permutations"$$T*+/(Q3;<8aq)8 < | q DL$8"49l;,,U3 --335 5~ &'nn A5  L q5AX&OVWw!nnT%:%:1%=%D%DEw W :  cJ7 99V  ( 3 3..x@   !"sQwq1ub"%A&,  # il "F&WW%%..u5 MTUW44ZCW UsGGsssGGGG#DI}= -67K-LM-LTQqD-LM E"&&u-C{   %  11':  11Hgg((2G8?8H8H8JK8J1^^A.8JMKs4m4sss4m444M8JQ8N#'S =#AIyy#tyy'::33ItIOJ!&'C!G !F1q5"b)+1%a(*Q-'*3<eT T++E2I"IIo6MM9;MM!Y%%7 )Y-=-=c-B-J-J-L-L$(M"YT Tr(c V^8dQhRRRR/#)rr?rr!rr#)r$s"r%r&rsszr(c J\V4# \\3dR#i;i)z0Convert value to int, returning None on failure.N)rr7r8)r?s&r%rjPallasKernel._safe_ints' s8O:&  s ""c V^8dQhRRRR/#)rshaperr!r#)r$s"r%r&rsY9r(c \V4p^.V,p\V^, RR4F(pW#^,,W^,,,W#&K* V#)z0Return C-contiguous strides for the given shape.r)r2rn)rrrrs& r%rm"PallasKernel._c_contiguous_stridessL J#'q1ub"%A Q%A,6GJ&r(c$V^8dQhRRRRRR/#)rrrrr!rr#)r$s"r%r&rs"I FVr(c /p\V4Fwr4WB9dR#W2V&K .pVF,pVPV4pVfR#VPV4K. \\ V44\V48wdR#V#)zMap 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)rrr r2r)rr stride_to_dimr}rXmappingr0s&& r%r PallasKernel._map_coeffs_to_dimss)+ g&DA! ! ' A!!!$Ay NN1   z'" #s6{ 2r(c V^8dQhRRRR/#)rctxrr!ztuple[bool, bool]r#)r$s"r%r&rs&&/&>O&r(c RpVPP4Fip\PP V4pVfRpK*VP 4F+pVP V4pV^8Xd RV3uu#VeK)RpK- Kk RV3#)z9Return whether an output has a zero or unknown dimension.FT)rvaluesrrtry_get_bufferrrj)r/r%has_unknown_dimrrr|dim_ints&& r%_zero_dim_output_flags#PallasKernel._zero_dim_output_flagss00779H''((2C{"&||~..-a<00?&*O & :o%%r(cV^8dQhRR/#)rr!tuple[int, ...]r#)r$s"r%r&rsdd_dr(c aaaaaVP'gR #VPP4UUu.uFwrVP'gKVNK pppVPP4UUu.uFwrVP'dKVNK pppV'd V'gR #VP 4pV'dV^8:dR #VPP4EF4wrgVP V4pVfKVwrrp \ V 4oS^8dK7V U u.uFqPV 4NK p p \;QJdRV 4F 'gK RM RM !RV 44'dK\\\,V 4oTpVFpVPV^4pK \4o\PP!V4F[pVP"'dKVP%4wpp VPV4pVfKAV^8gKJSP'V4K] S'gEK]\)VV3Rl\+S444pV'gEK\ S4^8dc\ S4\ V48XdG\ V4^8d7\,;QJd.V3RlV4FNK 5u#!V3RlV44u#EK\/\1S44oV^, S,o\2;QJd/V3Rl\+S^, 44F 'dK RM$ RM !V3Rl\+S^, 444pV'dPVR ,pTpV^8d!SSV^, ,8d V^,pK'\\+VV^,44pM3\)VVV3Rl\+S444pV'g \V4p\,;QJd.V3RlV4FNK 5u#!V3RlV44u# \/\1VPP5444pVP7V^,4pV^8wdVPV4M^oSf^oVP7V^,4pV^8wdVPV4M^pVf^pSV8dR #R #uuppiuuppiuup i) a8Determine 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. c3("TFqRJxK R#5irfr#rbs& r%rd3PallasKernel._get_reduction_axes..s6o9orTFc3H<"TFpSV,S9gKVxK R#5irfr#)rcrr_coeffsrs& r%rdr1s@Iqx)?Is" "c34<"TF qS, xK R#5irfr#rcrnds& r%rdr1 s 9ARc3X<"TFpSV,SV^,,8xK! R#5irgr#)rcrrs& r%rdr1s"S] WQU^ ;]s'*c3v<"TF.pSSV,8:gKSV,SS,8gK*VxK0 R#5irfr#)rcrr_stridespanrs& r%rdr1s;&71:-29!*th2NA&s 99 9c34<"TF qS, xK R#5irfr#r5s& r%rdr1"s.ARr7r)rr)rrHrqr_compute_reduction_numelrir2rjrkrrlrsubsrr%Add make_args is_number as_coeff_Muladdsortedrnrr3r4ror'rz) r/rrr_varspw_vars red_numelr load_indexr^rhrsactual_stridesrXstrides_or_noner_onlypvtermrzr0matchedrinnerstartaxesr_coeffpw_coeff pw_strider6r3r:r;rs & @@@@@r%_get_reduction_axes PallasKernel._get_reduction_axess$$$L $ 5 5 ; ; =P =!! =P!%!6!6!U!>ann11!>UWL113 INL$($9$9$?$?$A H((2D|04 -AAXBAv:HI.Q~~a0.OIs6o6sss6o666!%d3i!AG FR+(2 H ++F3>>>,,.qNN5)=QULLO 4@E"I@G 8}q x=CL0S[1_ 5 9 9595 9 999DN+HMX-DCSU2PQ6]SCCCSU2PQ6]SSM aiD7519+=$=QJEE%34"2Y =D5..5 .5... .C%BH$t44;;=>? ""6!9-.5l4>>'*  H##GAJ/08A DNN8,1  I i K qQU JsS S 8SS$Sc V^8dQhRRRR/#)rprefixesrr!rr#)r$s"r%r&r2s  j Z r(c ^pVFFpW0P9gKVPVPV,4pVfR#W$,pKH V#)zBCompute total numel for given prefixes (e.g., pointwise prefixes).N)numelsrj)r/rYr>r1numels&& r%_compute_prefix_numel"PallasKernel._compute_prefix_numel2sGAKKt{{1~6=   r(cV^8dQhRR/#)rr!rr#)r$s"r%r&r=s  * r(c ^pVPFApVP'gKVPVP4pVfR#W,pKC V#)zCompute total reduction numel.N)r.rrjr\)r/r>treer\s& r%r>%PallasKernel._compute_reduction_numel=sK$$D   tzz2= %  r(cV^8dQhRR/#rr!rr#)r$s"r%r&rHsCCtCr(c 2aSP'dR#SP4pVe V^8dR#.pSPPFpSP V4pVfR#VwrVrxp V 'gR#\ VRR4!4p V e^RIp WP8XdR#\;QJd.V3RlV4FNK 5M!V3RlV44p RV 9dR#VPV 4K V'd\\V44^8dR#V'dw^p V^,F pW,p K ^pSPFBpVP'dKSPVP4pVfR#VV,pKD W8wdR#R#)aV Check if TMA (Tensor Memory Accelerator) approach can be used. TMA works for simple element-wise ops but not for: - Reductions (need different accumulation patterns) TODO: TMA supports float64 for loading but not for reductions - Broadcasting (inputs have different shapes or output differs) - Non-contiguous tensors (strided, transposed) FN get_dtypecR#rfr#r#r(r%r4PallasKernel._can_use_tma_approach..hsdr(c3F<"TFpSPV4xK R#5irfrjrcrXr/s& r%rd5PallasKernel._can_use_tma_approach..psD8aq 1 18s!T)rr>r. input_buffersrirrfloat64rr r2rr.rrjr\)r/reduction_numel input_shapesrr^buf_objrs buf_numelrJr buf_dtyper shape_tuple input_numelrXrrar\sf r%_can_use_tma_approach"PallasKernel._can_use_tma_approachHsk  & & &779  &?Q+>%' II++D((.D|JN GGy-   lCEI$ -  %D8D%%D8DDK{"    ,-,2 C < 89A= K!!__ %L((((( NN4::6E}$ E)L )*r(c V^8dQhRRRR/#)rrr"r!z'tuple[Any, Any, Any, list, bool] | Noner#)r$s"r%r&rs(K(KS(K-T(Kr(c \PPV4pVfR#VP4p^pVF"pVP V4pYFeTMT,pK$ .pRp\ VRR4!4p V 'd\ V RR4MRp V e\ \V44F+p VP W,4p VPV 4K- \V4^8XdV^,eV^,^8wdRpMo\V4^8d`^p \ \V4^, RR4F<p W{,p V eW8wdRpVP W;,4pVfK4W,p K> W#WGV3#)z~Get buffer metadata (buf_obj, buf_size, buf_numel, actual_strides, is_contiguous). Returns None if the buffer doesn't exist. NTrcR#rfr#r#r(r%r/PallasKernel._get_buffer_info..rr(r;Fr) rrrrrjrrnr2r )r/rrqrsrrrXsvalrJrr buf_strider actual_strideexpected_stridedim_sizes&& r%riPallasKernel._get_buffer_infos] ''$$T* ?##% A>>!$D !1q 8I " , =?8>WVXt4D  !3x=) $z} = %%m4* 8}!!!$0^A5F!5K$)MX""#s8}q0"b9A$2$5M$, 0P(- #~~hk:H+'3 :)]JJr(c V^8dQhRRRR/#)rr"r r!ztuple[int, OrderedSet]r#)r$s"r%r&rs''' 'r(c VPV4p.pVFZpW@P9gKVPV,pVPVP4pVfKIVP V4K\ ^pVF pWx,pK Wr3#)zNCompute expected output numel and used vars from iteration variables in index.)r*rHrjrrr ) r/r"r9used_range_lengthsr4rrrls && r%rp-PallasKernel._compute_output_numel_from_indexs,,U3 C+++--c2!^^ELL9 )&--j9  #A  L$&&r(c$V^8dQhRRRRRR/#)rr"r r9rr!r#)r$s"r%r&rs$,6 r(c \4pVF]p\P!W4p\P!WT4pVf^pVP V4pTP VeTMT4K_ V#)zD Extract coefficients of iteration variables from index expression. )rrr5r6rjrD)r/r"r9 coefficientsr4r:r;coefs&&& r%_get_index_coefficients$PallasKernel._get_index_coefficientssj $.< C*GGSH(@@OF~>>&)D   T%5T6 B r(c ,V^8dQhRRRRRRRRRR/#) rrsrlrJrrrrr!r#)r$s"r%r&rs<#### # ! #  #r(c ^.p\V4^8d`^p.p\\V4^, RR4F:pVP^V4VPW,4pVfK2Wh,pK< V'd \ V4p VF p W9gK R# R#\ RV44p VF p W9gK R# R#)zB Check if access pattern requires gather (non-standard striding). Tc30"TF qfKVxK R#5irfr#rbs& r%rd5PallasKernel._check_gather_pattern..s*Vn11ns Fr)r2rninsertrjr) r/rsrJrrexpected_stridesrrrexpected_stride_setractual_stride_sets &&&&& r%_check_gather_pattern"PallasKernel._check_gather_patterns3 x=1 O! 3x=1,b"5 ''?;>>(+6'#/O 6 ",-="> $2% !+*Vn*V V $0%r(c(V^8dQhRRRRRRRR/#rrr"r"r indexingrr!r#)r$s"r%r&rs28888" 8  8r(c $aVPR8wgVP'dV#SPV4pVfV#VwrVrxp SPV4wrSP 4p SP W+4p SP WhW4p\V3RlV44p\V 4\V 48;'d;\V 4^8;'d%V^8;'d\V 4\V48p\PP4PR8HpV '*;'d5\;QJdRV4F 'dK RM RM !RV44p\;QJdRV 4F 'gK RM RM !RV 44pT;'dV'*;'dWz8HpV ^8dSWz8wgV'g V'd=\V 4^8d-V'g%V'g\SP!V4RR7#V#) zVCheck if buffer access needs strided indexing due to size mismatch or gather patterns.r$c3X<"TFpSPV4^8wgK^xK! R#5irgrjrks& r%rd7PallasKernel._needs_strided_indexing..s# OHqq8IQ8NHs* *rc3("TFqRJxK R#5irfr#rbs& r%rdr$s< #1aTM>rFTc3b"TF%p\V\\,4'*xK' R#5irf)rrrrs& r%rdr's!U 1Jq#+$> > > s-/r)rrrirprBrrsumr2rrrrrorkrrD)r/rr"rr^rqrsrrrJrrr9 all_iter_varsrhas_non_unit_stridesbuf_effective_dimsnot_all_vars_usedris_known_non_contiguoushas_symbolic_coefskip_for_non_contiguoussf&&& r%_needs_strided_indexing$PallasKernel._needs_strided_indexings    &(*@*@*@O$$T* <OFJC9m"&"G"G"N ++- 33EE  $99 m  ! OH OO  NS/ / / /I" / /"Q& / /IX. 446;;uD&3"3# # < #1< < #1< 9  CU UCCCU UU # P PF P Py7P 1 *.?CWI"+%"66u=T r(c(V^8dQhRRRRRRRR/#rr#)r$s"r%r&r:s24444" 4  4r(c VP'gVPR8XdV#\PP V4pVfV#VP 4p\ V4^8Xd\RVPR7#\ V4^8Xd=V^,^8Xd/VPV4'g\RVPR7#\ V4^8d^VPV4pV'g\VPRR7#RVP9d\VPV4RR7#VP'd.RVP9d\VPV4RR7#V#)zP Adjust index expression based on buffer shape (0-dim scalar, multi-dim, etc.). r$rTr) rrrrrrr2rrRrDr)r/rr"rrqrsrs&&&& r%_adjust_index_for_buffer_shape+PallasKernel._adjust_index_for_buffer_shape:sX  ! ! !X%7%75%@O''$$T* ?O##% x=A "x/E/E  MQ  q ,,U33"x/E/E  x=1  44U;M &&00+++&"::5AQU ;;;48#5#55"66u=T r(c(V^8dQhRRRRRRRR/#rr#)r$s"r%r&rps>HIHIHIHI" HI  HIr(c 8VP'gV#\PPV4pVfV#VP 4p\ V4pV^8dV#VP V4p\ V4^8wdV#\\V44p\P!W(4p VP\P!W44p V eV ^8:dV#\PPPW), 4p \V 4p T ^8gY8dT#TPTR,4p T eY,^8wdT#TP$P'T4pTfT#TPTP(4p^pTF%pTPT4pTfTu#TT,pK' TeY,T8wdT#RT^, ,pT ^8Xd T RT 2pM T T RT 2p\+TRR7# \ \"3dTu#i;i)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. z:, rFrr)rrrrrr2r*r3r4rr5rjr6r0r1rr7r8rHrrrr)r/rr"rrqrsrvr9r4r:r;rr<last_dimr var_lengthrrrXr}r slice_strs&&&& r%_try_multidim_slice PallasKernel._try_multidim_sliceps%%%O''$$T* ?O##%8} !8O,,U3 y>Q O4 ?#&CCEO  7 7 F  >Vq[O!!**5+;< VJ >Z1O>>(2,/  x0A5O %%))#. =O^^ELL1  Aq!Ay NI   !4 !AO$(# ?!("VH-I!(:,b9I%HHA:& O s HHHc$V^8dQhRRRRRR/#)r load_exprr"rr.r!r#)r$s"r%r&rs!666?6s6r(c RV RV R2#)zGenerate 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. zpallas_permute(r+r,r#)rrs&&r%_gather_permute_expr!PallasKernel._gather_permute_exprs! 2dV155r(c V^8dQhRRRR/#)rvar_namer"r!rZr#)r$s"r%r&rscjr(c WP9dVPV,#VPPFp\V4P 4pVP V R24'gK8\ P!RVPR^4^,4FpVPV4pVfKVuu# K R#)zTrace 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. r \btmp\d+\bN) rr0_linesr"rrWrefindallsplit_trace_to_load_source)r/rlineline_strrefr>s&& r%r"PallasKernel._trace_to_load_sources )) )%%h/ /LL''D4y'')H&&(3'788zz-q1I!1LM33C8%!MN (r(c(V^8dQhRRRRRRRR/#)rrr"rr"r r!z_IndirectAccessInfo | Noner#)r$s"r%r&rs2I I I !I *4I #I r(c  aaaSPV4pVfR#VwrVpVUu.uFpSPV4NK pp\V4^8g;\;QJdRV4F 'gK RM RM !RV44'dR#\ \ \ ,V4p SPV4p \V 4^8wdR#V ^,p SPW;4p V ^8Xg\V \ 4'gR#SPV 4p SPV .V 4pVfR#V^,p\V 4pV\^V^, 48dR#\V 4pSPV4pVfR#SPP!V4pVeSPV4pVeVwppp\V4^8wdR#SPV^,4fR#\"P$!VV3RlV44p\"P$!VV3RlSP'V444pVV8dR#\)VV\+V 4VVVR7#uupi)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. Nc3("TFqRJxK R#5irfr#rbs& r%rd7PallasKernel._detect_indirect_access..s'Hd333#'..1F1Fs1K1R1R"SS F>sA/A A)rrrrrr)rirjr2rkrrlrrYrrrmrmaxr"rrrrrr*rr)r/rrr"buf_inforhrsrX buf_size_raw buf_size_ints indirect_varsrrzrr"rrvindirect_var_namerindices_graph_name indices_info indices_size indices_numel iter_productrrrsf&&& @@r%_detect_indirect_access$PallasKernel._detect_indirect_accesss;((.  'Q13;<8aq)8 < | q CC'H<'HCCC'H<'H$H$H#'S <#@ //6 }  "$Q' ++E@ A:Zs33,,];**E7G< ?qz =! 3q$(+ + -223DE  "66::=I  )001CDL'+7(<Aq|$)>>,q/2: $ *+*!  $yy)#77>)  !L0"m,%*'   k=sIcV^8dQhRR/#)rr!rr#)r$s"r%r&r! s( ( t( r(c \4pVPFBwr#\P!RV4F"pVP VP 44K$ KD .pVP PFp\V4P4p\P!RV\P4pV'd5VPVP ^4VP ^4V34K~VPRWv34K RpV'dRp\V4FnwrpV 'gKW9gK\P!RV 4F;pVP 4V9gKVP VP 44RpK= Kp KVU UUu.uFwrpV e W9gKVNK uppp VP nR#uuppp i)aBRemove 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. rz^(tmp\d+)\s*=\s*(.*)NTF)rrrfinditerrDgroupr0rr"rmatchDOTALLr r) r/ live_varsrh store_linem assignmentsrrchangedrrhss & r%_eliminate_dead_indirect_code*PallasKernel._eliminate_dead_indirect_code! sg&0\ !33MA[[ ; aggi(<4 :< LL''D4y'')H0(BIIFA""AGGAJ D#AB""D(#9: (G$,[$9 q8 5[[<779I5%MM!'')4&*G=%:&1 %0!T8#8 D%0   s G /G c ,V^8dQhRRRRRRRRRR/#) rrr"rr"r rrr!r#)r$s"r%r&rK s<99 99 9 " 9 9r(c a VP'dVPWV4pVe5VPeWPP8wd \R4hWPnV R2#RVnVP P V4VP\P\P\\4pVP'dRMRpV'dRVP RV R2M VPpV R V R 2#V R VP R 2p VPR 8XdVP'gVP!W#4p V e"VP#W4p WP$V&V #VP'W#4p V eUV wo p R V RV R2p S VP(V&\*;QJd.V 3RlV 4FNK 5M!V 3RlV 44VnV #)z3 Build the load expression based on indexing mode. z-only one indirect access per kernel supportedz[0]Tz jnp.int32z jnp.int64rXrr,z[...].flatten()[rrr$zjnp.permute_dims(r+c36<"TFpSV,xK R#5irfr#)rcr1collapsed_shapes& r%rd0PallasKernel._build_load_expr.. s<8=1OA..)rrrAssertionErrorr rrDr/r%r rTorchMinTorchMaxrrrrrrrrrr)r/rrr"rindirect has_minmax idx_dtyperrr collapsedcpermrs&&&&& @r%_build_load_exprPallasKernel._build_load_exprK s  ! ! !33CuEH#''3 #7#77,K(0$c{"(,D %  ( ( , ,T 2599eii8LJ'+{{{ IH&&'y 1=''  U*3%q1 1%q!3!3 4A6I!!U*4;;;66tC# $ 9 9) JI8<//5 !% D DT QI ,1:.&7 {"UG1$M ?N55d;6;e<8=<ee<8=<73 r(c$V^8dQhRRRRRR/#)rrr"rr!r#)r$s"r%r&r s!sssr(c aVPR4'gV#\;QJd3V3RlVPP4F 'gK RM( RM$!V3RlVPP44pV'dX\P P V4oSe5SP4p\V4^8XdVR,^8XdRV R2#V#)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,) rc3<"TFjpVPR4'*;'dH\PPV4;oRJ;'d\ SP 44^8HxKl R#5i)rN)rWrrrr2r)rcrrqs& r%rdBPallasKernel._maybe_squeeze_intermediate_buffer.. su 4##E* * - -GG..x88E - -G$$&'1, -4s!A5*A5%A5TFz jnp.squeeze(z , axis=-1)r) rWrkr.rmrrrrr2)r/rr has_1d_inputrsrqs&&& @r%"_maybe_squeeze_intermediate_buffer/PallasKernel._maybe_squeeze_intermediate_buffer s u%% s !II33  sss !II33    gg((.G""++-x=A%(2,!*;))J??r(c(V^8dQhRRRRRRRR/#)rrr"r"r rr!r#)r$s"r%r&r s8YKYKYK *YK7:YK YKr(c ,\PPV4pVe\VP 44^8wdV#VP R4'dV#VP VP 4^,4pVfV#\PPV4pVeVP'gV#RpVPPFp\PPV4p V fK(\V P 44^8gKHV P 4U u.uFqP V 4NK pp \;QJdRV4F 'dK RM RM !RV44'dMRpK Ve\V4^8:dV#VPV4p \V 4^8wdV#\\V 44p WP9dV#VPV ,p VP V P 4V8wdV#VPP#4UUu.uF=wrVP VP 4V8XgK'VP$'dK;VNK? ppp\V4^8wdV#Rp\'V4UU u.uFwpqV8XgKVNK ppp \V4^8Xd V^,pMN^pVPP#4F.wppVV 8XdTpM!VP$'dK%V^, pK0 VfV#V\V4^, 8XdV#^.\V4,pR VV&V RRP)\+\,V44 R2#uup iuuppiuup pi) aReshape 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. Nrc3("TFqRJxK R#5irfr#rbs& r%rd:PallasKernel._maybe_broadcast_1d_buffer.. s;l}lrFTrr+r,r)rrrr2rrWrjrfis_floating_pointr.rmror*r3r4rHrrrqrrr'mapr")r/rr"rrq buf_lengthr  ref_buf_sizer other_bufrXr9used_varrrr matching_varsaxis_posr matching_dimspw_idxsym reshape_dimss&&&& r%_maybe_broadcast_1d_buffer'PallasKernel._maybe_broadcast_1d_buffer s0''$$T* ?c'"2"2"45:  ??5 ! ! ^^G$4$4$6q$9:   !!$'  U%<%<%<   //H**84I$Y-?-?-A)BQ)F;D;M;M;OP;Oaq 1;O P3;l;333;l;;;# 0  3|#4#9 ,,U3 y>Q  Y( 00 0 %%h/ >>%,, ': 5 --335 5~~ahh':5 >?nn A5  }  " '0'>R'>tq!z/'> R }  "$Q'HF//557Q(?%H~~~aKF 8    s<(1, , sS.. !# XIdiiC0F&G%HJJo Q. Ss$2N$N )N =N ( N8Nc$V^8dQhRRRRRR/#)rr"r rrr!r#)r$s"r%r&r s$>>>+:> >r(c VPR8wgVP'dV#VPV4pVP4pVP V4p\ VR4'd VP M \4V,pWe, pV'd\V4^8:dV#RpVPP4FwrVP V 4p V 'gK W8wdK(VPV 4p \ V R4'd V P M \4V,p W8wgWe8XdKvVPWV 4'dKRpM V'd\VPV4RR7#V#)z Check for im2col-like patterns where store uses block variables but load doesn't. For cat/expand patterns, both load and store prepared indices share block vars. For im2col patterns, store compresses to block vars but load doesn't. r$rAFTr)rrr+rBr*hasattrrArr2rrq_check_load_is_strided_inputrrD)r/r"rr)rCstore_orig_varsstore_prep_varsnew_varshas_im2col_patternrrIload_orig_vars prep_loadload_prep_varss&&& r%_check_im2col_pattern"PallasKernel._check_im2col_pattern sy    &(*@*@*@O..u5'') 2259~~66  ' '   #43/14O#$($9$9$?$?$A H!55jAN!0--j9I9n55&&\ N/?3U44n&*"5%B8 "66~F"  r(c(V^8dQhRRRRRRRR/#) rrr"rIr rrr!rr#)r$s"r%r&r; s,&9&9&9)3&9EO&9 &9r(c ~\PPV4pVfR#\VRR4!4pVfR#\VRR4pVfR#VP 4p.pVF]p \ P !W)4p \ P!W4p V fK5VPV 4p TPV eT MT 4K_ \4p \V4FQwrVPW~,4pVe V^8gK)VPV4pT PVeTMT4KS \V4V 8H#)zL Check if load coefficients match buffer strides (strided input vs im2col). NFrcR#rfr#r#r(r%r;PallasKernel._check_load_is_strided_input..E sDr(r;) rrrrrrr5r6rjr rrrD)r/rrIrrr buf_strides buf_sizes load_coeffsr4r:rint_coefbuf_stride_setrrXrint_ss&&&& r%r)PallasKernel._check_load_is_strided_input; s  gg  * ;lL9; >fh5  LLN  !C*GG XH&>>xMD>>$/""x/C8N "$k*DA~~il3H8a<q)""E,=51E + +&.88r(c V^8dQhRRRR/#)rrr"r!rr#)r$s"r%r&rc s2222r(c dVP'dR#VPV4pVfR#Vwr4r5p\V4^8wg\V4^8wdR#VPV^,4pVPV^,4pV^,pV^,p Ve!V eW8dVeVeV^8dV^8gR#VPP FZp VPV 4p V fKV wr.- sIjNHxjsrc3&"TFpRxK R#5i)rNr#)rcrhs& r%rdrV2 s7JqsJs) rrrDrnr2r r'rrrrrHr)r/r6r7rLrr@rNrHrGrrrO index_partsr| index_tuplerQrRr output_ndimnum_iter_vars_in_storetotal_kernel_iter_varsremaining_dimsis_element_wiserr n_leading n_trailing leading_onestrailing_nonesindirect_reshapedrhs&&&&&& r%_build_scatter_store_expr&PallasKernel._build_scatter_store_expr s(++,>F ""3'Vn #l2U '7L'7L'7LKS./&&&|4&&s+ 0 ))K0KU(;-y R |STUZT[[\] ]$N3 "=1 !,/ gg  &-0_c#,,.)! !$[!1C O!C!$T%:%:!;$q #n 4 A A&*@@  :EF+$8+KFK(IZJ1}a')3 !)J!6'3nAl^3~FVVW$X!$0!   0 1   IjI I)44 13 K4   | ,   7J7 7ii , e8K= +b AeWTU V +G 5s ;I> Ic$V^8dQhRRRRRR/#)rrr"r"r r!rr#)r$s"r%r&r: s!---Z-K-r(c VPPV4p\PP V4pW P V&VP V4pVPWV4pVPW!4pVe!W`PV&VPW64pM6VPWV4pVPWV4pVPW1W%4pVP'g5VPR8Xd$VP!W4pVP#WV4pVP$P'VP(VVR7pVPPV4p WP*\-V4&WP.V &V#)Nr$r )r.rrrrfrrrrr rrrrrrrrr.r/r0rr"r) r/rr"rr rrrcse_var buf_params &&& r%loadPallasKernel.load9 sliiood#!!$'',d#''.//XF //<  /5 & &t ,//<SPSV\R4R7#)inf)default)rr)r4r"r/s&r%_coeff3PallasKernel._handle_mixed_indexing.._coeff s..uc5<.P Pr(Tkeyr jnp.arange(r,Fz, 1z .reshape(-11r+z ).reshape(r4rcV^,#r=r#rqs&r%r5PallasKernel._handle_mixed_indexing.. s!A$r(c38<"TFqS8gK ^xK R#5irgr#rcr0 var_coeffs& r%rd6PallasKernel._handle_mixed_indexing.. s%NAI aa c38<"TFqS8gK ^xK R#5i)rNr#r{s& r%rdr} s*7!y=AA7r~rTz[:rc38<"TFqS8gK ^xK R#5irgr#rcr0indirect_coeffs& r%rdr} sI{!.6HAA{r~c38<"TFqS8gK ^xK R#5irgr#rs& r%rdr} sJ 1>7IQQ r~rSrr$z...]z[...)r*r(r)r2r-rEr,rYr"rHrrrreplacerBrr'r sortrr')&r/r"used_iter_vars_setrrr(r4 iter_coeffsrindirect_var_symsr rrXindirect_coeffsris_reduction_var range_entry range_size renamed_size arange_exprpaired_indirectunused_iter_vars n_output_dimsr trailing_ones reshape_exprr shape_parts shape_strall_componentsn_trailing_itern_trailing_indirectr` trailing_dimsr_ leading_nonesrbrr|s&ff @@r%r#PallasKernel._handle_mixed_indexingi s=&"55e< ""#56 ! "a '::e$ $  Q 2M.<=nsvc{n =JJt33E:;  33E:->?->cS-> ?7HH6G3q66!9,6GH ~ ! #M(:a(? #C3xHt,,,XX1F1Fs1K1X1X  $///"&"7"7"%?z)TUV !* 1 1(K HI#4&  !C  ! !6#;"< ="$C  ! !6#; C"@ A%= /FAs3xHd+++"33C8 (// #33J? "3K +DJJ|,D+EQG #&%N%N"N&)*.557*'#-/BB >$,z$9M%0MM?!"DK%--h D 50:*L,\:NI{IIIJ JJJ1}a (9 4 !)J!6".qs>BRRST Q (9 4 ".qtD a!)J!6".tN3C1E + !)), EI+*.}> @Is*T1T  Tc ,V^8dQhRRRRRRRRR R /#) rrr"r"r r7rr@rr!rr#)r$s"r%r&r, s4:7:7:7 *:73>:7FI:7 :7r(c `VeVR8wd\RV R24hVPPV4pVPP V4\ P PV4pVRJ;'d\VP44^8HpV'dRV R2V RV RV R2.pMVPeVPWSR 4pMVPW!4p V e@VPPVPV44VP!WSWV4.pMGVP#V4p VP%WV 4p VP'W*4p VP)WQW#W4pVF:p VP*P-V 4VP.P1W[34K< R#) NrBzpallas store mode 'z' not supportedr:r,r;z2.shape, _val) if _val.ndim == 0 else _val.reshape(rEF)rhr.outputstore_buffer_namesrDrrrr2rrr=_detect_scatter_patternr(r)r*rdrrrrIstoresrrr ) r/rr"r7r@r6r is_scalar store_linesrLrrs &&&&& r%storePallasKernel.store+ s   4 3D6IJ Jiit$ ##D)gg  &tO@@CLLN(;q(@ %eWA.%(-_`c_ddklK**6"??ER  $;;EH +''..t/G/G/NO66 D#K $33E:H $BBX H $99%JH#'"8"85#K D KK ! !$ '  " " ) )3+ 6 r(c(V^8dQhRRRRRRRR/#)rr"r r4rrqz int | floatr!r#)r$s"r%r&ri s,    , 7B   r(c VPV4pV^8Xd\P!W4p\V4# \\ 3dTu#i;i)z=Get integer coefficient of a variable in an index expression.)rzr%diffrr7r8)r"r4rqrzs&&& r%r#PallasKernel._get_index_coefficienth sO  C  A:JJu*E u: :& N s ;AAc$V^8dQhRRRRRR/#)rr"r output_namer"r!dict[str, Any] | Noner#)r$s"r%r&ru s)NNN.1N Nr(c ,VPV4p\V4^8wdR#V^,p\V4p\VP W44pV^8XdR#VP V4'gVP W%V4#VPWV4#)zDDetect scatter operation pattern. Returns scatter info dict or None.N)rYr2r"rrrR_detect_point_scatter_detect_iter_scatter)r/r"r indirect_syms indirect_symrrs&&& r%r$PallasKernel._detect_scatter_patternu s//6 }  "$Q' <( !$"="=e"RS Q ''..--kX X((nMMr(c(V^8dQhRRRRRRRR/#)rrr"rrrr!rr#)r$s"r%r&r s,   .1 CF  r(c  V'gR#\PPV4pVP4Uu.uFp\ V4NK pp\ T4^8dR#^p\ T4^, p\\ T4^, R R 4Fp Y78XdT pMYvT ,,pK RTRTR.R.RRRT/#uupi \ dR#i;i) z&Detect single-element scatter pattern.NrrrQrRrNTrOr)rrrrr Exceptionr2rn) r/rrrrrXrO cumulativerr|s &&&& r%r"PallasKernel._detect_point_scatter s ''$$[1C,/LLN;NqCFNL; | q  <(1, \*Q.B7C+"  s+ +J 8 L L 2 "  L   !<  s"1C CC C CCc(V^8dQhRRRRRRRR/#) rr"r rr"rrr!rr#)r$s"r%r&r s,+ + + /2+ DG+ + r(c aVPV4p.pVFp\VPW44pV^8gK&W`P9gK8VP VPV,P 4pVfR#VP \V4Wx34K VP SVR 34VPRRR7\V3Rl\V44R4p V fR#^p \WY^,R4FwrpWz8wdR#W,p K W:8wdR#RSRT RVRV U U U u.uF wrqV 3NK up p p R WY^,RU U U u.uF wrqV 3NK up p p R R R R/#uup p p iuup p p i)z0Detect scatter pattern with iteration variables.NcV^,#)r9r#rqs&r%r3PallasKernel._detect_iter_scatter.. sAaDr(Trtc3F<"TFwpwpq2S8XgKVxK R#5irfr#)rcrrrhrs& r%rd4PallasKernel._detect_iter_scatter.. s$ R&9?1ltQ\=QQQ&9s! !rrrQrRrNFrOr) r*rrrHrjrrr r"rr3rr)r/r"rrr(all_varsr4rzrr indirect_posexpectedrhrrs&&f& r%r!PallasKernel._detect_iter_scatter s11%802!C33E?@EqyS$9$99(=(=c(B(I(IJ>S5 9: " ~r:; .$ 7 Ri&9 R     (2B2D)E F Af   H!G  % L L (=L2IJ2IwqQF2IJ :J:L1MN1MgaAq61MN  D   KNs 6E2E9c ,V^8dQhRRRRRRRRRR/#) rr r r reduction_typerQr7z%CSEVariable | tuple[CSEVariable, ...]r!r#)r$s"r%r&r sCnnnn& n 5 n / nr(c aSP'g \R4hVR8XdSPW4#\V\4'd \ R4hW#V3pVSP P9dSP PV,#RRRRRR R R R R RRRR/p\.R"O4p\;QJdV3RlV4F 'gK RM RM!V3RlV44pSPV4p SP4p \RSPP444p T;'d)V RJ;'dV ^8;'dT ;'dV ^8p T;'dV ^8;'dV RJp VR8XdEV 'd6SP4p\!V4^8Xd V^,MTpRV RV R2pMRV R2pMVR#9d=Wc,pV 'd$SP4pV RV RVR$, R2pMV RV R2pMW69dbWc,pV 'd8SP4p\!V4^8Xd V^,MTpV RV RV R2pMCV 'd V RV R2pM2V RV R2pM)\ RV R\#VP%44 R 24hSP P'SP(VVR!7pVSP PV&V#)%a Generate code for reduction operations in JAX/Pallas. Reductions in Pallas work by: 1. Loading the input data into the kernel 2. Applying JAX reduction operations (jnp.sum, jnp.max, etc.) 3. Storing the reduced result The reduction happens over the loaded block of data. z!expected to be inside a reductionwelford_reducezHTuple reductions (e.g., welford_combine) not supported in Pallas backendrzjnp.sumrzjnp.prodrzjnp.maxminzjnp.minrkzjnp.anyargmaxz jnp.argmaxargminz jnp.argminc3@<"TFqSP9xK R#5irf)r[)rcr1r/s& r%rd)PallasKernel.reduction.. sI6H,6HsTFc3L"TFqP'gK^xK R#5irgr)rcrs& r%rdr s 9%=O=OAA9$ $Nxor_sumzjnp.bitwise_xor.reduce(z, axis=r,rXz, keepdims=True)z , axis=0)zReduction type 'z8' not yet supported in Pallas backend. Supported types: z , xor_sumrh)rmrz)rrr)inside_reductionrwelford_reduce_fallbackrrrhr.reduction_cacherrkr]r>rrHr'rVr2rlrIr/r0)r/r r rr7 cache_key reduction_opspointwise_prefixes has_pointwisepointwise_numelron_reduction_dimsis_partial_reductionis_symbolic_partialrR axis_exprreduction_expr reduction_opr>sf&&&& r% reductionPallasKernel.reduction s7"$$$ !DE E - -//= = eU # #Z  6 00 088++I6 6 9 J 9 9 9 l l (8I6HII6HII &*&@&@AS&T&*&C&C&E 00779    % %t+ % %!# % %  % %!1$   N N.2 N N$7N  Y &#//1'*4yA~DG4 #:5' ST!U#:5'!C 3 3(8L#//1$0>5'b !!L$0>5'!;  ,(8L#//1'*4yA~DG4 #nAeWGI;>NO%%1>5'!C$0>5'!;">"23$$(););)=$>#?yJ  "" LL # /5  + r(c V^8dQhRRRR/#)r buffer_namer"r!rr#)r$s"r%r&rI s&&3&4&r(c \PPV4pVP4pVP 4#rf)rrrrr)rrrs& r%_buffer_is_contiguous"PallasKernel._buffer_is_contiguousH s1gg  -!##%%r(cV^8dQhRR/#rdr#)r$s"r%r&rN sDD4Dr(c P aaa VP'dR#VP'dR#VP'dR#VP'deRP RVP P 44pVPF(pW P9dK\V4V9gK'R# \VPPP44pRpVFpVPV4pVfKVwrxryp\V 4^8gK2\V4^8gKDV ^,p V ^,p VPV^,4p VPV^,4p V fKV fKW8gKV fKV fKV ^8gKV ^8gKRpM \!VP"4;'gTVnVP'4pVRJ;'dV^8pV'd\)RVP*44M^Vn.o VFpVPV4pVfR#VwrxpVUu.uFpVPV4NK upo\.;QJdRS4F 'gK RM RM !RS44'dR#\S4\S 48gKSo K S 'gR#VP0e\VP04o \S 4p\VPP24V,pRpVEF.pVPV4pVfR#Vwrxp\V4^8XdK5WPP49d\VP4V,4oMVP0eWS9d\VP04oM^VUu.uFpVPV4NK upo\.;QJdRS4F 'gK RM RM !RS44'dR#\S4pVV8XEddWPP69pRp\9V4F@pSV,S V,8Xg%SV,^8XgS V,^8Xg V'dK>RpM V'dWPP"9d~VP"V,o\S4V8XdW\:;QJd*VVV 3R l\9V44F 'dK RM RM!VVV 3R l\9V444'gR#M V'dR#V'gU\.;QJd)VV 3R l\9V44F 'gK RM RM!VV 3R l\9V444'dRpEK]EK`VV8dRp\9VV, ^,4FSpRp\9V4F3pS V,^8XdKSVV,,S V,8wgK1RpM V'gKQRpM V'gR#RpEK\=\?S4\?S 44F#wppVV8wgKV^8wgKV^8wgK!R# EK1 V'gR#\@PBPE4PFR 8HpV'g ^@Vn$R#RVn$R#uupiuupi) aCheck 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 c38"TFp\V4xK R#5irfr"rcrs& r%rd1PallasKernel._can_tile_cpu_tpu..g s$O;N4SYY;NNTc3L"TFqP'gK^xK R#5irgr)rcras& r%rdr sD.d2C2C.rc3("TFqRJxK R#5irfr#rbs& r%rdr s/h9hrc3("TFqRJxK R#5irfr#rbs& r%rdr s3(QDy(rc3<"TFNpSSV,,SV,8H;'g(SSV,,^8H;'g SV,^8HxKP R#5irgr#)rcrint_sizer ref_shapes& r%rdr sd &3%T!W-1=11'Q0A511(|q01&3s%AAAc3r<"TF,pSV,SV,8H;'d SV,^8xK. R#5irgr#)rcrrrs& r%rdr s9%*QK9Q</DDIaL14DD*s77r)%rr rr(r'r0rr r"rlr.output_buffersrIrir2rjrrtile_has_transposer>rr.tile_skip_last_nrkrrmrr rnroziprrrrrr)!r/ compute_textrout_bufshas_col_major_outrr^rhrsrJr,r-d0d1ro has_reductionrXref_ndall_bufs has_tileablebuf_nd is_stridedmismatchrfoundrokrrrrrrs!& @@@r%_can_tile_cpu_tpuPallasKernel._can_tile_cpu_tpuN s ;;;  $ $ $  ! ! !    99$O4<<;N;N$OOL..:::w<</ / 005578 " H((2D|04 -AA>"a'CMQ,>#A&#A&^^HQK0^^HQK0NQQ(,%)!*#'t'B'B"C"X"XGX 779't3KK!8K  Dd..D D  "  H((2D|#' AAq3;<8aq)8,Vn VV,Vn!VPE4Vn#Rp VPH'dRp VPKVV4Vn&\OVPP4p!V!'dVPS4\4p"V"P34;_uu_4VP6P8FpV"P;\V44K RRR4VPUVVV"4VPW4#RV R R!PYVPB4 V R"2p#VP;V#4VP34;_uu_4VP[VVV4RRR4VP;R4V R#2p$.p%^\'VP<4,p&\+VP@4F>wppVV9gVP R4'gK&V%P]VV&,4K@ V%'d(R$R!PYR%V%44,R&,p'MR'p'\#\%^\'VP<4,44p(R$R!PYR(V(44,R&,p)VP;R)V) R*V' R+24R,R-.VP<,VP@,p*VP;RV$ R$R!PYV*4 R"24R!PYR.VPL44p+VP_V4wp,p-RMp.VP34;_uu_4V,'dVPaV.4EMV-'dDVP;R/4VP34;_uu_4VPaV.4RRR4VP;R04VPbe VP;R1VPb R&24VPLF3wp/p0VP@V/,pVP;V R2V R3V0 R424K5 VP;R54VP;R64VP;R74VP;R+4VPF'dVPeV4EM4VPgVVP@4VP@FYpVPiV4p1V1'dVPjPmV14MRp2V2fK?VP;V R2V R8V2 R+24K[ VP;R94VP;R:4VP;R74VP;R+4VP;R;4VP;R<4VP;R=R!PYVP@4,R>,4VP;R+4VPH'dXVPF'd$VP;R?4VP;R@4M"VP;RA4VP;RB4.p3VP<Fp4V3P]V4 RCV4 24K VPH'd#V3P]RD4V3P]RE4V3'dRFV RGR!PYV34 RH2p5MV RI2p5VPn;'d%VPp;'dVPs4p6V6'dVPuVV54MCVPn'dVPwVV54MVPyVV5VPLV+4RRR4VP{VV$4VPW4#uupiuupiuupiuupiuupp iuupiuupiuuppiuuppi +'giELo;iuupiuupiuupi +'giEL;i +'giELp;i +'giELX;i +'giL;i)Na} 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 in_out_ptrz2Pallas backend requires at least one output buffer cpur=r>TrFrrrrrrrrrrrrrrrrrrNr#z&, _pallas_tile=None, _pallas_ax2g=Nonedef _kernel(r+): _jit_wrapperrXc38"TFp\V4xK R#5irfrrcrms& r%rd.PallasKernel.codegen_kernel..s,L^SVV^r,)z()c38"TFp\V4xK R#5irfrrs& r%rdr0PAQr+@functools.partial(jax.jit, static_argnums=z, donate_argnums=r, out_shapes out_dtypesc34"TFwrV RV 2xK R#5i: Nr#rcros& r%rdrs%Pfq2aSkz+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 = (rz.reshape(_pallas_out_shapes[])zout_shapes_pallas = tuple(z& jax.ShapeDtypeStruct(shape, dtype)z; for shape, dtype in zip(_pallas_out_shapes, out_dtypes)rout_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 [rz_pallas_tile = _tilez_pallas_ax2g = _ax2gz$_pallas_tile = _pallas_out_shapes[0]z_pallas_ax2g = {}=z_pallas_tile=_pallas_tilez_pallas_ax2g=_pallas_ax2gfunctools.partial( _kernel, ),z_kernel,)r r )r in_ptrr#)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])>rr.python_argdefsrrWrr0r' RuntimeErrorrrqrr"rrrrrlrnr2rrraliasable_out_ptrs_codegen_importsindent_codegen_iteration_varsr0rrrrrrr tile_cpu_tpur _compute_alias_pairsrrrr _codegen_scalar_prefetch_wrappergetvaluer'_emit_kernel_bodyr r+ writelinesr_codegen_tiled_specsrrrrrrrv_codegen_jit_wrapper_tma_codegen_jit_wrapper_legacy_gpu_codegen_jit_wrapper_cpu_tpu_codegen_main_entry)7r/rrarg_defs call_argsrhrrr1rrsize_var_namesrouterrPrrrrrrrrrrflagrrrr% kernel_bodyrextra_kernel_paramsuse_scalar_prefetchkernel_body_spkernel_signaturejit_wrapper_namedonate_indices base_offsetdonate_literalstatic_argnumsstatic_argnums_literalwrapper_paramsalias_map_literal has_zero_dimr)zero_dim_return input_idxout_idxrcshape partial_argssv_param kernel_arguse_tmas7&& r%codegen_kernelPallasKernel.codegen_kernels %)II$<$<$>!Q)12A 2&3Om||I7N11mO$ $! 5N(OAA}  $DII$6$6$=$=$?@&3KmN7J11mKST T!% 8 8 > > @ @ %% EL @  --o 77>>@EEN&6FG+-$E%)OE "%+: *9_U=S=S ugV /  % $! 5M(NAA}  +\9)M9&$3$9$9$; H$;jdD4TT$; H  "&uS-?'@"A  [[["$ "+=!9#!9IC,,!9 #   # ;; .  0  (  , ( , "6 , & & !4  2 0! "!4# $% (#2 c"$&    ! !  ( (c : ++%%c$i0,"%)II$<$<$>!Q)12A 2#DII$6$6$=$=$?@*7O-Q;Nqq-O$ $! 5M(NAA} #/1A1A"A!- !=!224   ' ' '"J 33CI"4#7#78   . . 0+-N&&(( LL//D",,SY70)  1 1  ==? ""+htyyAWAW7X6YZmYnnpq '( [[]]  " "4c : r)],7#c1122 "3#:#:;IC $)F)F%%cK&78<  499,L^,L#LLtSN!NeAC,?,?(@$@AB!$tyy0P0P'P!PSW!W  ''=&>>Oa ! < (3+>+> >AXAX X  ./q>1J0K2NO II%P%PP(,(C(CC(H% o  [[]]0"NN#PQ8'E..:NN01L1L0MRP+.//&Iw33I>ENN 'UG+GyPRS+: ;<GHQs#$$$--c22249P9PQ!$!8!8#'#:#:5#A (!99==hG!% "- NNeWCwixq+QR"9NN#>?NN#PQNNUNN3'NN#=>NN#RSNN(499S5L5L+MMPSSNN3'///((('=>'=>'MN':; " # 3 3H ''8*AhZ(@A!4/// ''(CD ''(CD#5k])DIIVbLcKddf!gJ$/=!9JKK55..55224 11#zB[[[88jI55Z:Kqx   &67}}}3O L   I#>" ! !3O *)((]]T' ]snn$n0n nnnn"4n"n(* n(9n-n- n2 n2 1 n8n8A n>o oo(oo 8o!:o5#p?)p(p :CpA5pCp2pBp)p00p!p5p$p31p> o ! o2 5 p p  p p- c$V^8dQhRRRRRR/#)rr%rrrr!rr#)r$s"r%r&r(s$    />   r(c j.p\VP4Fwr4VPR4'dOVPVR4'd5V R2pVPP V4pVP Wc34KhKjVPP V4pVP Wc34K V#)r FrF)rrrWrrr"r )r%rrrOr alias_namerNs&& r%r0!PallasKernel._compute_alias_pairs's.0 &s'8'89MGy))"&&tU33$(6J # 7 7 = =j II&& ';<4  3399$? ""I#78:r(c(V^8dQhRRRRRRRR/#)rrrr?r%rr!rr#)r$s"r%r&r7s2+++$+ +  +r(c @VPFUp\V\4'd"VPVP 44K:VPP V4KW VP F(wrVWSP9gKVPV4K* R#)z:Emit the kernel body lines and store operations into code.N)rrr"rrr rr)r/rr?r%rr rs&&&& r%r3PallasKernel._emit_kernel_body7sr &&D$$$t{{}- ""4( ' $(#9#9 G000z*$:r(c(V^8dQhRRRRRRRR/#) rr%rrr"r?rr!rr#)r$s"r%r&rGs8m8m8 m8m8$ m8  m8r(c  VPf \R4hVPpVPp\VP4pVP Uu.uF1pWtP 8wgKWtP8wgK'Wv9gK/VNK3 ppVP .VP.,V,\VP4,VP,p VPRV RRPV 4 R24VP4;_uu_4VPWSV4RRR4VPR4V R2p R R .VP,VP ,p \\^\!VP4,44p R RPR V 44,R ,p VPRV R24VPRV R RPV 4 R24VP4;_uu_4VPpVP pVP"p\!VP$4pVPR4\V4F$pVV8wgK VPRV RV R24K& VPRV R24V^8XdVPRV RV R24MFV.\V4Uu.uFpVV8wgK VNK upO5pVPRV RV RV RV R2 4.pVFpRV R2pVPRV RV R24VPR V R!V R"24VPR#V R$24VPR V R!V R%V R&24VPR'4VPR V R!V R(24VP'V4K .pVPF1pRV R2pVPV R!V R"24VP'V4K3 VPUu.uF pV R)V 2NK ppV'dR*V R+RPV4 R2pMV R,2pVPR-4^\!V4,pVPR.4VP4;_uu_4VPR/4VFpVPR0V R1V R224K VPFpVPR34K RRR4VPR4\!VP4pVPR4RPR5.V,4,R,4.p\)VP4F,wpp^V,V,p VP'V R6V 24K. RPV4p!\V4Uu.uF pR7V R82NK p"pR9RPV"4,R,p#VPR:4VP4;_uu_4VPV R;24VPR4VPR?4VPR@4VPRA4RRR4VPRB4V'dVPRCV! RD24VP*'g VPREVP, R;24RRR4V.RF.,V,V,p$VPRGRPV$4 R24VPRH4RRR4VP/W4R#uupi +'giEL\;iuupiuupi +'giEL;iuupi +'giEL;i +'giL;i +'giL;i)Iz=Emit kernel, JIT wrapper, and main entry for scalar prefetch.Nz"expected indirect_access to be setrrr+rr#rrrrXc38"TFp\V4xK R#5irfrrs& r%rd@PallasKernel._codegen_scalar_prefetch_wrapper..rrrrrr,z_D = 1z _D = _D * z.shape[rz_seq = z .shape[0]z _table_3d = rz.shape[0], 1, _D)z .transposez ], 1, _D)rh_3dzif z.ndim == 2 and z.shape[0] == _seq: rz.reshape(_seq, 1, _D)zelif 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)))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_SPECrz/jax.ShapeDtypeStruct((_seq, 1, _D), out_dtypes[r!r_result = pl.pallas_call(,z out_shape=z'grid_spec=pltpu.PrefetchScalarGridSpec(znum_scalar_prefetch=1,z grid=(_seq,),zin_specs=_in_specs,zout_specs=_out_specs,r'zinput_output_aliases={  },z interpret= _table_3d)(z?return tuple(r.reshape(s) for r, s in zip(_result, out_shapes)))rrrrrrrrrlrrr'r-r3rrnr2rrr rrrr9)%r/r%rr?rr alias_setr1other_input_paramsprefetch_kernel_paramsrDrJrHrItableindicesind_dimrvrr}rpallas_call_other_argsp3dpallas_call_alias_argssvrQ kernel_refnum_non_alias_in_specsrh num_outputsalias_map_partsrOarg_idxrKout_shape_partsout_shape_exprall_pallas_argss%&&&& r%r1-PallasKernel._codegen_scalar_prefetch_wrapperGs    ' !EF F''xxs//0 ,, ,*** ))) " A,   # # $##$ %  !3##$ %    ;-x 2H(I'J" M [[]]  " "4c : r)],7 < (3+>+> >AXAX X eAC,?,?(@$@AB!$tyy0P0P'P!PSW!W 9:P9QQR S  ./q>1J0K2NO [[]]((E,,G++Gx++,D NN8 $4[<NNZwgaS#BC! NNWWIY7 8!|eWIeWDUVWKeDk"JkQ'\11k"JK"5'D6gWWIY8&( "'!CjQCqc9KLMcU#aS0EFGqc56cU#aSs*MNw'cU#aS0BCD&--c2(&( "%%!Cj#c!,ABC&--c2& 473F3FG3FRrd!B4L3FLG( 6 ,/03 !,}G4  NNA  &'-C)D%D " NN? +O2CNN"e$)),.UV2))ANN<0* NN3 c//0K NN 499k][-H#IICO  !O'(8(89 44w>&&'"WI'>?:!% / : {++AB!BG+ !499_#==CN NN6 7*Q/0N+;1=>HI[[]]NN#;<NN?3NN#89NN#:; # t$"NN%=>O=PPT#UV{{{NNZ0E0E/Fa#HI  - ()))  NNR / :;1= > NNQ qx   7A (]]>#K:H":#]] C]s]?, ]?>]?]?^A_"?A _" ^ -^ 3D_"^"A4_"A^" -C _"6^6A _"A_ "A^; ' _ &_ /_ A_" ^  _"" ^3 -_"; _  _  _ _"" _2 c V^8dQhRRRR/#rr%rr!rr#)r$s"r%r&rs--O--r(c BRpVP'dVR, pVR, pVR, pMVP'g VR, pVP'd%VP'gVR, pVR, pVPP VRR7R#) a 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')TstripN)rrrrsplice)r/r%importss&& r%r,PallasKernel._codegen_importss" ::: , ,G K KG R RG%%% R RG     T G K KG t,r(c V^8dQhRRRR/#)rrrr!rr#)r$s"r%r&rs,:r(c @^p^p\RVPP444pVPP4FQwrVWQ8Xd!VP'gVu#WC,u#VP'd V^, pKHV^, pKS R#)zMap 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. c3L"TFqP'dK^xK R#5irgr)rcrs& r%rd2PallasKernel._get_iter_var_axis..sS<NN11._get_nd_shape_if_matches..9s5'A%Qemm(<==A1D'sAA)NN)rrr(r2rrrr)rrrr\ iter_lengthss& r%_get_nd_shape_if_matchesRPallasKernel._get_reshape_target_shape_and_numel.._get_nd_shape_if_matches5s''((2C{c#,,.1Q6!!EEEEIIe$E%*l%:E> L Lr(N) rrHr'rrrrr%rrcopyrr.rm) r/rrcandidate_buf_namesreshape_target_shapereshape_target_numelrr>rs & @r%#_get_reshape_target_shape_and_numel0PallasKernel._get_reshape_target_shape_and_numel(s "..557 7Aahhemm(<=AHH 7   M#77<<>""499#:#:;5942+H-h7Fayy=C:$#99 , $999 s 2C-C-c$V^8dQhRRRRRR/#)rbroadcast_varslist[_BroadcastedIterVar] broadcast_idxrr!r"r#)r$s"r%r&rLs$7HK r(c W,pVPPpVPV4pVPV4p\ V4pVP WV4pR.V,p WiV&RP V 4p RV R2p V RV R2p V #)rwr+rvr,r)rrrr,r-r2_broadcast_axis_idxr') r/rrbvrrrenamed_length length_strnum_broadcast_dimsaxis_idxrrarangereshapeds &&& r%$_make_broadcasted_iteration_var_expr1PallasKernel._make_broadcasted_iteration_var_exprLs *--f5ZZ/  0++ += e00 *HIIk* zl!,XYyk3r(c$V^8dQhRRRRRR/#)rr?rr%rr!rr#)r$s"r%r&r_s.eTeT)eT0?eT eTr(c  4aVP'd%VP'gVP'gR#VPR4VP 4wr4\ VPP 44p.pRp\V4FMwowrVPV P4p V e W8XdSpK1VP\SWV 44KO \V4p \V4EFwowrWP9dK\V4p V Pp VPV 4pVPV4pVPV 4p V f}V'd\V ^8dUSV8wdN\!V3Rl\V44R4pVe+VP#VV4pVPV RV 24KVPV RV R24KV'dR\V4^8dBW8Xd<RP%RV44pRV R2pVPV RV R V R24EK>V ^8dQSV8wdJ\!V3R l\V444pVP#VV4pVPV RV 24EKR P%R VP&P(44pV V9pT;'d]VP'*;'dDVP*'*;'d+VP,'*;'dV P.'*pV'dVP1V4MRpVe<VPV R V RV RV R24VP2P5V4EKVPV RV R24EK R#)Nz*# Define iteration variables as JAX arraysc3R<"TFwrVPS8XgKVxK R#5irfrrcrrrs& r%rd7PallasKernel._codegen_iteration_vars..s!Q'@tqAEESL'@' 'rz = jnp.arange(r,r+c38"TFp\V4xK R#5irfrrbs& r%rdrs%K6Jc!ff6Jrrvrc3R<"TFwrVPS8XgKVxK R#5irfrrs& r%rdrs#%";$!quu|AA";rrc38"TFp\V4xK R#5irfrrs& r%rdrs(S?RtT?Rrz = jnp.arange(_pallas_tile[z$]) + pl.program_id(_pallas_ax2g.get(z, 0)) * _pallas_tile[r)rHrr(rrrlrqrrjrrr rr2r"r,r-r3rr'r0rrr rrr rD)r/r?r%rr var_itemsr total_var_idxrrrrrrrrrrrrrrvar_in_computecan_tile_relativerrs&&& @r%r.$PallasKernel._codegen_iteration_vars_s %%%dkkkd>Q>Q>Q JK  4 4 6 3..4467  %.y%9 !C!' 5J%**L # %%'WZH &:!0%.y%9 !C!'1117|H\\F!11&9NN3J/J!(*Q.},$(Qy'@Q%M%0#HH*M $-- #dV.DE %% . A&NO%,-16 II%K6J%KK &zl!4%% #fXYykQR&ST#a'C=,@ $%"+N";%! @@"M%% #dV&<= $yy(St||?R?R(SS !)\!9"// KK// 222//!555//"... "9JD++G4t'))#*$?zJ==EJG++3*A7 0044W=))XJnZLPQ*RSQ&:r(c(V^8dQhRRRRRRRR/#)rrrrrrr!r#)r$s"r%r&rs066166 6 6r(c \;QJdRV4F 'gK RM RM !RV44p\;QJdRV4F 'gK RM RM !RV44pT;'dTpV'dV#V^, V, #)c3j"TF)p\VP4PR4xK+ R#5irNr"rrWrcrs& r%rd3PallasKernel._broadcast_axis_idx..s)! 6DC O & &s + +ns13TFc3t"TF.p\VP4PR4'*xK0 R#5irrrs& r%rdrs,! :HBBJJ**3/ / /.s68)rk)rrrhas_reduction_varshas_pointwise_varsis_mixeds&&& r%r PallasKernel._broadcast_axis_idxs!S! 6D! SSS! 6D!  !S! :H! SSS! :H!  &<<*<  !A% 55r(c$V^8dQhRRRRRR/#rr%rrSr"r!rr#)r$s"r%r&rs,vUvUOvUvUQUvUr(c n VPpVPpVPpVPR4VPR4VPR4VPR4VPR4VFpVPRV R24K VPR4VPR4VPR 4VUu.uFqw R 2NK ppVUu.uFqw R 2NK p pVUu.uFqw R 2NK p pVUu.uFqw R 2NK p pVPR 4VPR 4W,p W,p \ \ V44Uu.uFpRV 2NK ppRP W,4pVPRRP V 4 RV R24VP4;_uu_4VPR 4VPR4VP4;_uu_4VPR4VPR 4VPR4\\W44F$wpwppVPRV RV RV R24K& VPR 4VPR4\V4FwppVPRV R24K VPR 4VPR4W,pVPR4P4pVPV RRP V4 R24VPR 4VPR4VPR 4\W4FwppVPR!V RV R"24K VPR#4VPR 4VPR$4R%R%R%4VPR 4VPR&4VPR'4R%R%R%4VPR 4VPR(4VPR)4\V 4F&wppWN,pVPR*V R+V R,24K( \V 4FwppVPR*V R-V R.24K VFpVPR*V R/24K VPR 4VPR04VPR14VPR 4VPR24VPR34VP4;_uu_4VPR44VPR54VPR64R%R%R%4VPR74VFpVPR8V R924K VPR4VPR 4VPR:4VPR;4R%#uupiuupiuupiuupiuupi +'giELG;i +'giEL ;i +'giL;i)?@nEUUW X [[]] NN2  NN: ;EFr"VW-6)=.)A)NN27);`ah`iituvtwwxy.r"AB%&78DAqNN%A!A#FG9r"89#4#I &--c288: )Adii8H.I-J!LMr"M45*-.@*U&HhNN28*BxjHlm+V;<r"}-EH NN2  NN5 6 NNK LSX r  O  -.&'89MAz,/J NN#J * NN7 8 NN< = t(E NNT% 4 5) s r => STGGAFAT]]P]sO W#" W(5 W- W2W7;XF>W< >;X24X$< X  X X! $ X4 c$V^8dQhRRRRRR/#rr#)r$s"r%r&rRs)L L "L 03L L r(c VPpVPpRRPV4 R2pVPR4VPR4VFpVPRV R24K VPR4VPR4VPR 4VPR 4VPR 4VPR 4VPR 4VPRV R24VPR4VPR4VPRV,4VPR4VPR4VPR4VPR4VPR4VPR4VPR4VPR4VPR4VPRV R24VPR4VPR4VPR4VPR 4VPR!4VPR"V,4VPR#4VPR$4VPR%4VPR&4VPR'4VPR(4VPR)4VPR*V 24VPR+4VPR,4VPR-4VPR!4VPR"V,4VPR#4VPR$4VPR.4R/#)0rr+rz7# Check if all tensors have same size (no broadcasting)z_all_sizes = []z_all_sizes.append(rrz' _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)r#z if _can_pad:z5 # All tensors same size - safe to flatten and padz+ _padded_inputs = pallas_gpu_pad_inputs(r,zZ _aligned_out_specs, _is_scalar = pallas_gpu_align_output_specs(out_shapes, out_dtypes)z _result = plgpu.kernel(z z% out_shape=_aligned_out_specs,z )(*_padded_inputs)zD return pallas_gpu_unpad_results(_result, out_shapes, _is_scalar)rdzA # Different sizes - check if it's a reduction (scalar output)z) _out_numel = math.prod(out_shapes[0])rcz 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)N)rrr'r)r/r%rSrr input_listrs&&& r%r7,PallasKernel._codegen_jit_wrapper_legacy_gpuRsxx!55#678: PQ ()(E NN/wf= >) 12 @A ?@  V  r ~& NO DZLPQRS  h  45 zJ./ >? /0  R  w  O  BC v 01  U  H TUVW <= GH TU {# 89 ~ 23 BC 34 /0 {#  ^  >? 12 LZL Y  {# UV  l  89 ~ 23 BC 34  V r(c V^8dQhRRRR/#)rrr"r!rZr#)r$s"r%r&rs r(c tVPPP4Fwr#W18XgK Vu# R#)z:Map a kernel parameter name back to its graph buffer name.N)r.rmrq)r/r graph_name inner_names&& r%rPallasKernel._param_to_buf_names4&*ii&=&=&C&C&E "J"!!'Fr(c V^8dQhRRRR/#rr#)r$s"r%r&rsG G G DG r(c  VPpVPpVP'dRMRpVP'dRMRp.pVPFHpVP V4pTP V'dVPPV4MR4KJ VPp V 'dRV 2MRp \;QJdRV4F 'gK RM RM !RV44'dR \V4 2MRp VPR V R V R V V V R 2 4VPR4VPR4VPR4VPR4VPR4VPR 4VPW!P4VPFYp VP V 4pV'dVPPV4MRp V fK?VPV RV RV R 24K[ RPVP4p\;QJdRV4F 'gK RM RM !RV44'dgRPRV44pVPRV R24VPR4VPRV R24VPR 4R#VPR4VPRV R24VPR 4R#)u)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=r#c3("TFqRJxK R#5irfr#rcr1s& r%rd4PallasKernel._codegen_tiled_specs..s4)QD=)rTFz, 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_shapesrrr+c3("TFqRJxK R#5irfr#rs& r%rdrs0i}irc38"TFp\V4xK R#5irf)rrs& r%rdrs!=9a$q''9rz_perm_flags = [rr#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 [)rrrrrrr rrrrkrrrrr')r/r%rskip_n has_transposeis_tpu_literal all_permsr1rmgpmgp_arg perms_argrrPr perm_lists&& r%r5!PallasKernel._codegen_tiled_specssxx&&"&"9"9"9w #&:::735 ((A..q1H   =E++//94 ) ((14'u-"s4)4sss4)444d9o. /  &(!(#$%k'!  %  )* 56 23  T  89 s &&t-D-DE,,E..u5HDLT2266x@RVF!%E7)F81EF -YYs667 30i03330i0 0 0 !=9!==I NN_YKq9 : NN5 6 NN%%/LA  NN3  NN5 6 NNYZdYeefg  NN3 r(c ,V^8dQhRRRRRRRRRR /#) rr%rrSr"rrrKr!rr#)r$s"r%r&rs<#)#) #)#)+ #)  #)  #)r(c JVPpVP'dRMRpVPR4VPRV,4VPR4VPR4VPR4VPRVP R24VPR V R24TPV'dR V R 2MR 4VPR 4VP'dDVPUu.uF pRV R2NK ppVPRRP V4 R24VPR4VPR4VPR4VPR4VPR4VPR4R#uupi)_gridz(1,)rgrcz out_shape=out_shapes_pallas,z out_specs=out_specs_pallas,z in_specs=in_specs_pallas,z interpret=rhz grid=z input_output_aliases={ riz input_output_aliases={},rkzpallas_ensure_nonzero_rank(r,r+zif isinstance(_result, tuple):zF _result = tuple(r.reshape(s) for r, s in zip(_result, out_shapes))rdz, _result = _result.reshape(out_shapes[0])zreturn _resultN)rr/rrrr') r/r%rSrrKr grid_exprr1 kernel_input_params_nonzero_ranks &&&&& r%r8)PallasKernel._codegen_jit_wrapper_cpu_tpusuxx#000Gf  23 v *+ 9: 89 67 (=(='>a@A 9+Q/0 ++<*=T B/ t  " " "r#_mainrrXr+, stream=None):z*jax.config.update('jax_enable_x64', False)jax.clear_caches()rz.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)) _placeholderout_shapes = (tuple(rErNtorch_dtype_to_jax_runtime(rout_dtypes = (rrzexported = jax.export.export(z, platforms=['tpu'])(r,z'_'.join(str(s) for s in z + 'x' + z''zkernel_key = 'z _out_' + z + '_in_' + z.if not tpu_torch_pallas.lookup_custom_kernel('z', kernel_key):ztry:z)tpu_torch_pallas.register_custom_kernel('zF', kernel_key, serialized_mlir_module=exported.mlir_module_serialized)zexcept TypeError:z/', kernel_key, exported.mlir_module_serialized)zinput_tensors = [rzoutput_shape_tensors = [z torch.empty(z.shape, dtype=z, device='tpu'),z.dtype, device='tpu'),c34"TFwrV RV 2xK R#5irr#rs& r%rd7PallasKernel._codegen_main_entry_tpu..s)TO&1QCr!+Or z_input_output_aliases = { z }c3<"TFwr\V4xK R#5irfr)rcr_os& r%rdrs.U_'1s1vv_sz_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)zif _results is not None:z7for _in_idx, _out_idx in _input_output_aliases.items():z0input_tensors[_in_idx].copy_(_results[_out_idx]))rrrr'rr-rlrrr rrrrrrfr rrrr)r/r%rDr main_namekernel_name_strall_input_params param_nameall_jax_input_namesrYptrr dtype_exprsrr wrapper_placeholder_argsshape_key_partsr1output_key_exprinput_key_partsinput_key_exprinput_tensor_names alias_map_strdonate_argnums_strs&&& r%r$PallasKernel._codegen_main_entry_tpusxx r'u- // 9+Qtyy)?)?@A Q [[]] NNG H NN/ 0 $C$4$45S=M=M8NN . !l#j\2$%@ N/ NNH I"$ !.. !l"F!l"Ej\QY[$**j\+FG / ''e?e>se8M$**cU,+?@ ( NN ))@Q@QR@QvdV73@QRST  &(K))3377='GG--h7E(#**+=e+DE ""%@g#NO* NN+dii .DDtK L)5l'C $ $ + +C,?,? @ $ + +,? @ NN/#$%II678; !O&&&&)B1#W'MN'6E   1$ !O,,&&)B1#W'MN-6E   1$  NN  1"#~.0  NN@@QQ`a v&[[]]NN+,-RS# 23[[]]NN+,-;<#""&c&6&6!7$s?O?O:P!P  NN.tyy9K/L.MQO P NN5 6--D"77;;DAH+ ! 1 1( ; , NN".tfN5)K[ \%NN&tfN4&@VW. NN3  $ )TCOO)T T !<]O3OP%)YY.US__.U%U"!45G4HJK;<56 NN6 "'()67 NN. /'()BC NN5 6M[[]]NN#UV# ]FSf#]]#]],<$#] ]]sD1^?3\E^?A^?A^?8*] "\" 81] )\6 ?] B^? B ] *^??B:^?9]2 1^?^ 1^?*^, 1^ ^, ^?^?" \3- ] 6 ] ] ]  ^? ]/ ) ^?2 ^ = ^? ^ ^? ^)$^, , ^< 7^?? _ c$V^8dQhRRRRRR/#rr#)r$s"r%r&rs)NN"N69N Nr(c  8aVPpVPR4VP R2pVPRV RRPVP4 R24VP 4;_uu_4VPR4VP 'dVPR4VPR4VP'dNVPR 4VPF,pVPVVVPVP R 7K. VPR 4VPF:pVPR 4'gKVPW6VPR R 7K< VPR4VPF:pVPR4'gKVPW6VPRR 7K< VPR4TPRRPVPUu.uF pRV R2NK up4,R,4.pVPFvpVPPV4p V e@\P P#V 4p V eVP%\'V 44KaVP%RV R24Kx VPRRPV4,R,4/oVPF pV R2SV&K VPF pV R2SV&K RR.p V P)VP*4V P)V3RlVP,44VPRV RRPV 4 R24VPR 4VP.'dOVPR!4VP.F-p VPV ,p VPV R"V R#24K/ RRR4R#uupi +'giR#;i)$r#r rrXr+r z)jax.config.update('jax_enable_x64', True)z>jax.config.update('jax_default_device', jax.devices('cpu')[0])r z*# Convert Torch -> JAX for donated outputs) contiguousz+# Convert Torch -> JAX for in-place tensorsr Fz!# Convert Torch -> JAX for inputsr(Tz-# Prepare output metadata from PyTorch tensorr rrErNrrr_jaxrrc36<"TFpSV,xK R#5irfr#)rcr arg_name_maps& r%rd;PallasKernel._codegen_main_entry_default.. s%/Ft T""/Frzres = r,zjax.block_until_ready(res)z9result_values = res if isinstance(res, tuple) else (res,)z'.copy_(torch.from_dlpack(result_values[z])))rrrr'rr-rr_emit_torch_to_jaxrrrWrrrrrrfr r rrrr)r/r%rDrrrYrrrrr wrapper_call_argsrrr*s&&& @r%r(PallasKernel._codegen_main_entry_defaultsxx r'u-  9+Qtyy)?)?@A Q [[]] NNF G###T NN/ 0KL"%"2"2J++" ##7 ,#3 NNH I''>>,//++Dszze+T( NN> ?''>>(++++Dszzd+S( NNJ K NN ))@Q@QR@QvdV73@QRST  &(K))3377='GG--h7E(#**+=e+DE ""%@g#NO* NN+dii .DDtK L+-L!.. .8\+> Z(/'''*e4L S!(".| <   $ $S%8%8 9  $ $%/2/F/F%  NNV$4#5QtyyAR7S6TTUV W NN7 8&&&O22C"005HNN#*$KC5PST3C]FSG]]s<7APBPAP0APPFP+APP P c ,V^8dQhRRRRRRRRRR /#) rrrrr"rrr'r!rr#)r$s"r%r&rs;WWW(+W59WJNW Wr(c PV'dRMRpVPV RV V R24R#)z.detach().contiguous()z .detach()z_jax = jax.dlpack.from_dlpack(r,N)r)rrrr'suffixs&&&$ r%r,PallasKernel._emit_torch_to_jaxs-.8)[ (#A(F8STUVr(c$V^8dQhRRRRRR/#)rrr"rz IRNode | Noner!rr#)r$s"r%r&r"s!'''='D'r(c B\PPpVPP 4wrEpVUu.uFqwP NK ppVU u.uFqP R4'gKV NK p p \\\V44p \VR/4p V U u.uF4p V PV R4'gKWPV 4,NK6 p p V RRPW,4 R2pVPV4R#uupiuup iuup i)z7Generate the Python code that calls this Pallas kernel.r r+Fz.run(r+r,N)rr wrapper_coder.r)rrWrlrr"rrr"r'r)r/rrwrapperr:r;rhrkernel_param_namesr1r call_arg_strs aliasablealias_call_args kernel_calls&&& r% call_kernelPallasKernel.call_kernel"s''&&$(II$<$<$>!Q.67hffh7&8T&8LLrvrirprrrrrrrrrrrrrrr3r=rIrdtyping_extensionsoverriderkrrrrrrrrrrUr0r3r1r,rrrr.rr6r7rr5r8r9rrr,r<rM __classcell__rs@r%rrs &I)5E &5.+`8!>FbH<8:5K7IV , ,(ST DhTrh,U0*& dL  CJ(KT'( #J8t4lHIT66$I V( T9v4YKv>@&9P2h27:rN `- -^@D:7 :7x  N* B+ Zn`&& DLQf  + m8^->&":H&eTN66&vUpL \G R#)JD nW`N`WW ''r(rcVa]tRtRt]t]RRl4tV3Rlt]t ]t RRlt Rt V;t #)PallasSchedulingi7c V^8dQhRRRR/#)rrz torch.devicer!zOrderedSet[BackendFeature]r#)r$s"r%r&PallasScheduling.__annotate__;sEE,E;UEr(c 6\\P.4#rf)rrREDUCE_TO_SINGLE_ELEMENT)clsrs&&r%get_backend_features%PallasScheduling.get_backend_features:s>BBCDDr(c <\SV`W4'gR#VP4'dVP4'd^RIHp/pVP P F.p\WS4'gKVPWEP&K0 VP P FMp\WS4'gKVPV9gK)WEP,VP8wgKLR# R#)F) MemoryDepT) rcan_fusertorch._inductor.dependenciesrNrreadsrr"r)r/node1node2rNreads1rrs&&& r%rOPallasScheduling.can_fuse@sw--     E$6$6$8$8 >F((..c--'*yyF88$/((..c--#((f2Dhh'3994$/r(c(V^8dQhRRRRRRRR/#)rsrc_coder" node_schedulezSequence[BaseSchedulerNode]r'rr!r#)r$s"r%r&rGXs2"""3" " "r(c \PPpWP9dVPV,#\P P 'd%\V\P P 4MRp\P!VPR44P4R,pVR8XdRV 2pMRV RV 2pWtPV&VPRV4p\4pVPRV: R24VPVR R 7VPR 4\!W$4wrV R V 2p VP#WxP%4V 4V#) r#zutf-8:NNfusedpallas_rhr zasync_compile.pallas(z, r'''Trz''')r)rrr5 src_to_kernelr tritondescriptive_namesrhashlibsha256encode hexdigestrrrrr define_kernelr2) r/rWrXr'r6 fused_name kernel_hashrcompile_wrapperoriginsdetailed_originsmetadata_comments &&&& r%rdPallasScheduling.define_kernelXsP ''&& ,, ,((2 2}}... "-1P1P Q  nnX__W%=>HHJ2N  #K=1K#JK*5h'##O[A(*!!$9+"OPxt4!!&)$7 $O!%Yb)9(:;k+C+C+EGWXr(r#)rHrIrJrKr kernel_type classmethodrKrOcan_fuse_verticalcan_fuse_horizontalrdrMrBrCs@r%rErE7s9KEE *!"""r(rE)F __future__r dataclassesr`rrr@typingrrrr%rtorch.utils._ordered_setrtorch.utils._sympy.functionsrrr rr r#r runtime.runtime_utilsr r&rr virtualizedrblock_analysisrcommonrrrrrsimdrrrrr-r?collections.abcrNrOirrP ops_handlerrQ schedulerrR MAIN_SUFFIX_logginggetArtifactLoggerrHr]rVr*rhrj dataclassrrrrrrEr#r(r%rs`" ++ / 6>/CBM<&& 2+- ..228]K>>.U,Ur  Kr  j    '''.Z>':Z>'z|C~Cr(