KgjD SSKrSSKrSSKJr SSKJr SSKJrJ r J r J r J r J r \(a/SQO/SQrSrS7S\R S\R S-S \R 4S jjr\R$"S S S .5\R&"S8\VVs/sHnSHn\R)0XS9PM M snnS/S.\ D6\R*S\R,S\R,S\R,S\R,S\R,S\R,4 Sj555r\R$"SSS .5\R&"S8\VVs/sHnSHn\R)0XS9PM M snnSS/S.\ D6\R*S\R,S\R,S\R,S\R,S\R,S\R,S\R,S\R,4Sj555r\R$"SSS S!.5\R&"S8\VVs/sHnSHn\R)0XS9PM M snnS/S.\ D6\R*S\R,S\R,S\R,S\R,S"\R,4 S#j555r\R$"S$S%S&S!.5\R&"S8\VVs/sHnSHn\R)0XS9PM M snnSS/S.\ D6\R*S\R,S\R,S\R,S\R,S\R,S\R,S\R,S"\R,4S'j555r\ S\R6S \4S(j5rS9S\R S\R S-S)\R S-S*\S+\R6S-S \R 4 S,jjrS:S-\R S.\S/\S0\R S-S\R S-S1\S2\S+\R6S-S \R 4S3jjr "S4S5\RBRD5r#\RHRJS9S\R S\R6S-S)\R S-S*\S+\R6S-4 S6jj5r&gs snnfs snnfs snnfs snnf);N)prepare_chunk_indices)IS_AMDIS_NPUautotune_cache_kwargsget_multiprocessor_count input_guard tensor_cache))r r r r ix cu_seqlensreturncBUbUR5S:XdS5eURup#nUS:XdS5e[R"U5nURSS- n[ U5HnXR 5nXS-R 5n X- n U S::aUSX24*USX24'ME[R"USX245n USXS- 24U SS&XSX24- n XSX24'M U$[R RS5n U "U5n X- n U $)NInput must be [B, T, D]*Batch size must be 1 when using cu_seqlensr)rrr)dimshapetorch zeros_likerangeitemnn ZeroPad2d)rrBTDresultNistartendseq_lenshifteddelta time_shifts ڥC:\Users\johnn\AppData\Local\Packages\PythonSoftwareFoundation.Python.3.13_qbz5n2kfra8p0\LocalCache\local-packages\Python313\site-packages\fla/modules/token_shift.pytoken_shift_refr-sIuuw!|666|''aAvCCCv!!!$   Q ! #qAM&&(EqS/&&(CkG!|()!UY,'7q%)|$ **1Q \?;5Q;/ AuyL/1',q%)|$ XX'' 6 Q-  cUSSL$Nrargss r,r4:d<0 a "--"2#Dy$*C77=372366rxx@ GGMC!G+a/ 0 3 3BHH =77:+,bggj36F6J.KS( 7R<s37S1WMS( 7R<+ (RYYq"% %C 'C 7E "7A#gga&js+O  -6C!GcMC!GcM ''%)"3#># zHWWQ\4FSLE VU-  37"1:a# a# ]2CcB/#r.cUSSL$r0r1r2s r,r4r4r5r.cUSSL$Ngrad_cache_outr1r2s r,r4r4d+;&>rxxH GGMGaK/!3 4 7 7 A77:+,bggj36F6J.KSB, 7R<s37S1WMSB, 7R<+ (RYYq"% %C 'C!*a# a# I 7E "7A#wwrF{- sQw WW]Y%>SI uz)uzH52772=s;;D fd-  3)3TD'#r.cf[R"UR55R5$N)rmaxdiffr)rs r,prepare_maxlensrws! 99Z__& ' , , ..r.r8 output_cacherlc^^^^URupVn[R"U5nTb[T5n[ T5S- mOUmUS:*n [ (aU (aTU-[ :aSn U(a-[R"TU4URURS9n OSn U (abTb[ T5S- mOUm[R"U5n TU4n US:H=(d US:H=(a UT:Hn [U "UUTUU UUU UU S9 O[S[R"[R"[SXV-5[!URR"5555nUcTb [%TU5nTb [ U5O[R"Xn5m[R"U5n [R"X{5m[R"XV-S5nUUUU4S jn [&U "UUTUUU UUU UUTUS 9 UTXiU 4$) NriFdevicedtype) rrPrr8rQr!r"r?rBrC@r c>TT-TbS4$T4$Nrr1metar$rgNTrs r,gridtoken_shift_fwd..gridRj.DLL!LLr.)r"r?rfrergrB)rr empty_likerlenr_NPU_MAX_TRITON_GRIDemptyrrtritonnext_power_of_2rbmincdivrrindexrry)rrr8rrlr r!r"rPuse_short_kernelrQr?rrCrfrer$rgrs ` @@@r,token_shift_fwdr|sggGA! A J '  Oa  Dy v"q1u/C'C KKAqxxqwwG    !J!#AA  # #A &1vF1qAv0!q& $T*!* V++FKKB F^_`_g_g_m_mFn,op q  Z%;1*bAM#-#9S v{{1?Q  # #A & [[  [[d #LL#D)     * ai //r.rr$r!dcacherhas_init_cachecd^^^^URSn[R"U5n [R"U5n U(a-[R "TU4UR URS9n OSn U(aTU4n [U "UU TUU UUU S9 X4$[S[R"[R"[SUR5U-5[UR R5555n UcTb [TU 5nTb [!U5O[R"X-5m[R"TURS-S5n[R"U5n [R"X5mUUUU4Sjn ["U "U UTUUU UUU U UTS 9 X4$) Nr r)rrrrr}r!r"r?rr rrc>TT-TbS4$T4$rr1rs r,rtoken_shift_bwd..gridrr.)r"r?rfrerg)rrrrrrrrrrrrnumelrrrrr)rr$r!rrrrrlr"r?rr}rrfrergrs ` ` @@r,token_shift_bwdrs  A    "B   " BaVBIIRXXN1v$T*! ) F  1V++FKKB a8P8PQSQZQZQ`Q`8a-cde  Z%;1*bAM#-#9S v{{1?Q [[RXXa[$ /  # #A & [[ LL#D)          r.c \rSrSr\\S S\RS\RS-S\RS-S\S\RS-4 Sjj55r \\SS \RS \RS-4S jj55r S r g) TokenShiftiNrrr8rrlcx[XX4U5upgpn X lXPlXplXlXlUSLUlXj4$r)rrrlr$r!r has_cache) ctxrrr8rrloutputr$r!rrQs r,forwardTokenShift.forwardsL 5DASXhu4v11 #)/T)   r.rrc [XRURX RURUR UR 5up4USUSS4$r)rr$r!rrrrl)rrrr grad_caches r,backwardTokenShift.backwardsK)UUCEE6>>),)=)=s}}cN_N_a4T4//r.r1NNFNr) __name__ __module__ __qualname____firstlineno__ staticmethodrrTensorbool LongTensorrr__static_attributes__r1r.r,rrsHLHM9= ! !%,,2E !||d* !AE !$//$6 ! !0%,,0 t0C00r.rcUb5UR5S:XdS5eURSS:XdS5e[RXX#U5upVU(aXV4$U$)a Token-shift operation implemented with Triton kernels. Args: x: Input tensor of shape [B, T, D] (or [1, T, D] when `cu_seqlens` is supplied). cu_seqlens: Optional cumulative sequence lengths of shape [B + 1]. When supplied, `x.shape[0]` must be 1 and `x.dim()` must be 3. cache: Optional cache tensor of shape [N, D] that holds the last token from the previous call. output_cache: Whether to return the updated cache alongside the output. In previous versions this parameter did not exist and the cache was always dropped; to preserve backward compatibility the default is False. Returns: output: Tensor of shape [B, T, D] after applying the token-shift. cache_out: Tensor of shape [B, 1, D] containing the last token that should be fed as `cache` in the next call. Only returned when `output_cache=True`. rrrrr)rrrapply)rrr8rrlrrQs r, token_shiftrsd:uuw!|666|wwqzQL LL"((][F   r.rr1r)NNTFN)'rrtriton.languagelanguagerI fla.ops.utilsr fla.utilsrrrrrr NUM_WARPS_AUTOTUNErrr- heuristicsautotuneConfigjit constexprrbryrrrintrrrrautogradFunctionrcompilerdisablerr<s00r,rsN /pp&,]2C '+! ||! t#! \\!H<?, +I#J  bI E# F+   K> ||K> K>||K>||K>||K>||K>  K>\<?, +I#J  bI E# F+  t :C ||:C :C :C :C :C||:C||:C||:C  :Cz<H@ , +I#J  bI E# F+   <A ||<A <A||<A||<A <A  <A~<H@ , +I#J  bI E# F+  t 7E ||7E 7E 7E 7E 7E||7E||7E 7E  7Et/ 0 0/S// '+!%-1 H0 ||H0 t#H0 <<$ H0 H0 ##d* H0  \\ H0^#'&*! -16 6 6 6 LL4  6  t# 6  66##d*6 \\6r0((00+/!%-1 $ ||$  4'$ <<$ $ $ ##d* $$w x X \ s $R9$R?$S $S