Kgj`SSKrSSKrSSKJr SSKJr SSKJr SSK J r SSK J r \R"SS05\"S#SVVs/sHnSHn\RS U0US 9PM M! snn/S QS .\ D6\R"S S/S9S\R S\R S\R S\R S\R S \R S\R 4Sj555r\"S5S$S\R$S\R$S\R$S\R$S\R$S\R$S\S\R(S-S\S \S!S4S"jj5rgs snnf)%N)dispatch)fla_cache_autotune)exp2)autotune_cache_kwargs IS_VARLENcUSSL$)N cu_seqlens)argss ڴC:\Users\johnn\AppData\Local\Packages\PythonSoftwareFoundation.Python.3.13_qbz5n2kfra8p0\LocalCache\local-packages\Python313\site-packages\fla/ops/kda/chunk_intra_token_parallel.pyr sd<0<)BH) num_warps)KHHV)configskeyTN)do_not_specializerrrBTBCc  [R"S5[R"S5nnU(aSnSUnn[S5HYnUU:dM UU-S-nU[R"UU-S-5R [R 5:aUnMTUS-nM[ Un[R"UU-5R [R 5[R"UU-S-5R [R 5nnUU- n UU- nO UU -U -nUU -nUU :agUU -nUU -U-nUU -nUUU--nX-nUUU -U -- nUUU -U -- nUUU -U -- nUUU -U -- nUUU -U-- nUUU -- n[ R"U 5n UU-[R"SU5-n!U!U-n"[R"SU 5n#U!U :n$U#U :n%U$SS2S4U%SSS24-n&U"SS2S4U -U#SSS24-n'[R"UUU -U --U'-U&SS9R [R5n([R"UUU -U --U'-U&SS9R [R5n)UUU -U --U!SS2S4U --U#SSS24-n*UUU --U!-n+[R"U*U&SS9R [R5n,U)[R"U+U$SS9R [R5SS2S4-n)[U[US-[U UU-555GHn-[R"UU-U -U --U'-U&SS9R [R5n.UU-U -U --U!SS2S4U --U#SSS24-n/[R"U/U&SS9R [R5n0[R"U%SSS24U.[U,U0- 5-S5n1[R"U(U1-SS9U-n2[R"U)U1-SS9[R"U-U:SS5-n3[R"UUU -U --U!U --U-U --U2R URR 5U$S 9 [R"UUU -U--U!U--U--U- U3R URR 5U$S 9 GM g) Nrrr)maskotherg)axisg?)r!)tl program_idrangeloadtoint32tritonnext_power_of_2arangefloat32minwherersumstoredtype element_ty)4qkgbetaAqkAkkscaler rrrrrrrrri_tgi_hgi_nleftright_midboseosi_ti_ci_si_tci_tsGBKo_hvo_ho_km_hvm_km_hkp_qkb_qb_kp_gp_betab_gjb_kjp_gjb_gjb_kgjb_Aqkb_Akks4 r )chunk_kda_fwd_kernel_intra_token_parallelr^s@q!2==#3$De rAe|e|)"''*s"2Q"67::288DDE7D 77:+,//9277:PSCSVWCW;X;[;[\^\d\d;eS #ISjqyAoQh ax )C 8 C 8D #(?DgAqNAqNAr!OA38B;C38B;CC"HD--a0B "9ryyB' 'D !)C ))Ar C "9D 'C 4=3tQw< 'D q$w'>DsU{+e3sU{+bhhq3wS.II sRx"}$tby01r69588CIIDXDX;Y`de sRx"}$tby014t;UXXciiFZFZ=[bfg:rkdar4r5gkr7r8r9r:r chunk_sizesub_chunk_sizereturnc ^^^/URQURSP7ummpmUb[U5S- OTn Un U nUUU4Sjn[U"UUUUUUUUU TU TU U US9 XE4$)a) Token-parallel implementation: each token gets its own thread block. Supports both fixed-length and variable-length sequences. Reduces wasted computation on padding. Writes directly to Aqk and Akk tensors (in-place). Args: q: [B, T, H, K] k: [B, T, H, K] gk: [B, T, HV, K] cumsum of gates (HV >= H for GVA) beta: [B, T, HV] Aqk: [B, T, HV, BT] output tensor to write to Akk: [B, T, HV, BC] output tensor for diagonal blocks (fp32) scale: attention scale chunk_size: BT (default 64) sub_chunk_size: BC (default 16) rrcB>TT-[R"TUS54$)Nr)r*cdiv)metaBrrs r grid0chunk_kda_fwd_intra_token_parallel..grids AE6;;r4:#>??r)r4r5r6r7r8r9r:r rrrrrrr)shapelenr^)r4r5r`r7r8r9r:r rarbrrrrrrirhrrs @@@r "chunk_kda_fwd_intra_token_parallelrm~s>+agg*rxx{*NAq!)5J!1A B B?-d3    " 8Orr )N@)torchr*triton.languagelanguager$fla.ops.backendsrfla.ops.utils.cacherfla.ops.utils.opr fla.utilsr heuristicsConfigjit constexprr^Tensorfloat LongTensorintrm)rrs00r rs %2!+< B%I  tRjI 6% 7   sCj)Zh ||Zh Zh ||Zh Zh Zh !Zh"||#Zh*Zhz %+/5 ||5 ||5  5 ,, 5  5  5 5  4'555 55O s&E