+ UV-j9.^RIHt^RIt^RIHt^RIHtHt^RIH t ^RI t ^RI HtHtHt^tRt^tRRlt]!RR 7R 4t]!RR 7R R l4t]!RR 7R 4t]!RR 7R4t]!RR 7R4t]!RR 7R4t]!RR 7R4t]!RR 7RRl4t]!RR 7RRl4t]!RR 7RRl4t]!RR 7RRl4t]!RR 7RRl4t ]!RR 7RRl4t!RRlt"R R!lt#R"R#lt$R$R%lt%R&R'lt&R(R)lt'RR*R+llt(RR,R-llt)RR.R/llt*]!RR 7R0R1l4t+]!RR 7RR2R3ll4t,]!RR 7R4R5l4t-]!RR 7R6R7l4t.]!RR 7RR8R9ll4t/]!RR 7RR:R;ll4t0RR<R=llt1RR>R?llt2RR@RAllt3]!RR 7RRBRCll4t4]!RR 7RRDREll4t5]!RR 7RF4t6RRGRHllt7RRIRJllt8]!RR 7RKRLl4t9]!RR 7RMRNl4t:RRORPllt;]!RR 7RRQRRll4t<]!RR 7RSRTl4t=]!RR 7RURVl4t>!RWRX]4t?!RYRZ]4t@!R[R\]4tA!R]R^]4tB!R_R`]4tCRaRbltDRRcRdlltEReRfltFRgRhltGRiRjltH]!RR 7RkRll4tIRmRnltJ]!RR 7RoRpl4tK]!RR 7RqRrl4tLRsRtltMRuRvltN]!RR 7RwRxl4tORyRzltP]!RR 7R{R|l4tQR}R~ltR]!RR 7RRl4tSRRltTRRltURRltVRtWRRltXRRltYRRltZRRlt[RRlt\RRlt]Rt^Rt_RRlt`RRltaRtbRRltc!RR4td!RR4te!RR4tf!RR4tgRRlth!RR4tiRRltj!RR4tk!RR]4tl!RR]4tmR#)) annotationsN) lru_cache) NamedTupleOptional) _BaseCachecreate_attention_maskcreate_causal_maskư>cV^8dQhRR/#)returnbool)formats"c/Volumes/fast/ai/experiments/ui-tars-smoke/.venv/lib/python3.14/site-packages/mlx_vlm/turboquant.py __annotate__rs<<$<cn\\R4;'d\PP4#)metal)hasattrmxr is_availablerrr_metal_availablers$ 2w  ; ;BHH$9$9$;;r)maxsizecv\4'gR#Rp\PPR.ROR.VR7#)Na auto lane = thread_position_in_grid.x; auto repeat_idx = thread_position_in_grid.y; auto n = thread_position_in_grid.z; auto token_count = norms_shape[2]; auto kv_heads = norms_shape[1]; auto repeat_count = q_rot_shape[2]; if (repeat_idx >= repeat_count) { return; } auto b = n / (kv_heads * token_count); auto rem = n % (kv_heads * token_count); auto h = rem / token_count; auto t = rem % token_count; auto q_ptr = q_rot + ((b * kv_heads + h) * repeat_count + repeat_idx) * Dim; auto packed_ptr = packed + ((b * kv_heads + h) * token_count + t) * PackedWidth; float acc = 0.0f; for (int d = lane; d < Dim; d += 32) { int bit_offset = d * Bits; int word_idx = bit_offset / 32; int offset = bit_offset % 32; uint value = packed_ptr[word_idx] >> offset; int spill = offset + Bits - 32; if (spill > 0) { value |= packed_ptr[word_idx + 1] << (Bits - spill); } value &= ((1u << Bits) - 1u); acc += static_cast(q_ptr[d]) * codebook[value]; } acc = simd_sum(acc); if (thread_index_in_simdgroup == 0) { out[((b * kv_heads + h) * repeat_count + repeat_idx) * token_count + t] = acc * static_cast(norms[(b * kv_heads + h) * token_count + t]); } turboquant_mse_scoreoutname input_names output_namessourceq_rotnormspackedcodebookrrfast metal_kernelr"s r_mse_score_kernelr,sB   'FP 77   #<W  rcV^8dQhRR/#r repeat_countintr)rs"rrrJs[[#[rc |\4'dV^8dR#.ROp\V4FpVRV R2RRV RV R2., pK V.RO, p\V4FpVR V R 2., pK V.RO, p\V4FpVR V R V R2., pK VRR., p\V4FpVRV RV R2., pK VR., p\V4FpVRV RV R2., pK VRR ., p\PP RV 2.ROR.RP V4R7#)uTiled MSE score: preload R queries into registers, loop over tokens. Processes TokTileSize tokens per threadgroup instead of 1. Reduces threadgroup count by TokTileSize×, amortizing dispatch overhead.Nz float qr_z[DimsPerLane];@ for (int i = 0, d = lane; i < DimsPerLane; i++, d += 32)z qr_z,[i] = (d < Dim) ? static_cast(q_base[z * Dim + d]) : 0.0f; } float ps_ = 0.0f; ps_z += qr_z [i] * code; } float s_ = simd_sum(ps_);z1 if (thread_index_in_simdgroup == 0) {( out[(bh * RepeatCount + z) * token_count + t] = kn * s_;turboquant_mse_score_tiled_r r). auto lane = thread_position_in_grid.x;+ auto n = thread_position_in_grid.z;r2* auto token_count = norms_shape[2];' auto kv_heads = norms_shape[1];K auto num_tok_tiles = (token_count + TokTileSize - 1) / TokTileSize;$ auto bh = n / num_tok_tiles;* auto tok_tile = n % num_tok_tiles;- int t_start = tok_tile * TokTileSize;A int t_end = min(t_start + TokTileSize, (int)token_count);r2z7 auto q_base = q_rot + (bh * RepeatCount) * Dim;: auto pk = packed + bh * token_count * PackedWidth;+ auto nm = norms + bh * token_count;r2zK // Preload query values into registers (constant across all tokens))r2z3 // Precompute bit offsets for key unpackingz6 int k_words[DimsPerLane], k_offs[DimsPerLane];z# bool k_spills[DimsPerLane];B for (int i = 0, d = lane; i < DimsPerLane; i++, d += 32) {z int bo = d * Bits;z! k_words[i] = bo / 32;z k_offs[i] = bo % 32;z0 k_spills[i] = (bo % 32 + Bits) > 32;r4z0 constexpr uint mask = (1u << Bits) - 1u;r2zB // Token loop: key data changes, queries stay in registers/ for (int t = t_start; t < t_end; t++) {+ auto pt = pk + t * PackedWidth;z1 float kn = static_cast(nm[t]);r2)z3 for (int i = 0; i < DimsPerLane; i++) {z8 uint vv = (pt[k_words[i]] >> k_offs[i]);zX if (k_spills[i]) vv |= pt[k_words[i]+1] << (Bits - (k_offs[i]+Bits-32));z1 float code = codebook[vv & mask];r#rrangerr)r*join)r/linesrs& r_mse_score_tiled_kernelrSIs   !1 E&<  s. 1NaS LQCOc d  ! E(<  )!H566! E<  's'!K@AA!  E <  (?1#R@AA! ;E<  6qc9WXYWZZ[ \  ! E 77  *<. 9<Wyy  rct\4'gR#Rp\PPRR.R.VR7#)Na auto word = thread_position_in_grid.x; auto row = thread_position_in_grid.y; if (row >= values_shape[0] || word >= PackedWidth) { return; } auto values_ptr = values + row * Length; uint packed_word = 0u; int start = max(0, (int(word) * 32 - (Bits - 1)) / Bits); int end = min(Length, ((int(word) + 1) * 32 + (Bits - 1)) / Bits); for (int idx = start; idx < end; ++idx) { int bit_offset = idx * Bits; int word_idx = bit_offset / 32; int offset = bit_offset % 32; uint value = values_ptr[idx] & ((1u << Bits) - 1u); if (word_idx == word) { packed_word |= value << offset; } if (word_idx + 1 == word) { int spill = offset + Bits - 32; if (spill > 0) { packed_word |= value >> (Bits - spill); } } } out[row * PackedWidth + word] = packed_word; turboquant_pack_lowbitvaluesrrr(r+s r_pack_lowbit_kernelrWsC   F> 77   %JW  rct\4'gR#Rp\PPRR.R.VR7#)Naa auto idx = thread_position_in_grid.x; auto row = thread_position_in_grid.y; if (row >= packed_shape[0] || idx >= Length) { return; } auto packed_ptr = packed + row * PackedWidth; int bit_offset = idx * Bits; int word_idx = bit_offset / 32; int offset = bit_offset % 32; uint value = packed_ptr[word_idx] >> offset; int spill = offset + Bits - 32; if (spill > 0) { value |= packed_ptr[word_idx + 1] << (Bits - spill); } out[row * Length + idx] = value & ((1u << Bits) - 1u); turboquant_unpack_lowbitr&rrr(r+s r_unpack_lowbit_kernelrZsC   F& 77   'JW  rcv\4'gR#Rp\PPR.ROR.VR7#)Nau auto lane = thread_position_in_grid.x; auto repeat_idx = thread_position_in_grid.y; auto n = thread_position_in_grid.z; auto token_count = norms_shape[2]; auto kv_heads = norms_shape[1]; auto repeat_count = q_proj_shape[2]; if (repeat_idx >= repeat_count) { return; } auto b = n / (kv_heads * token_count); auto rem = n % (kv_heads * token_count); auto h = rem / token_count; auto t = rem % token_count; auto q_ptr = q_proj + ((b * kv_heads + h) * repeat_count + repeat_idx) * Dim; auto packed_ptr = signs + ((b * kv_heads + h) * token_count + t) * PackedWidth; float acc = 0.0f; for (int d = lane; d < Dim; d += 32) { int word_idx = d / 32; int offset = d % 32; uint bit = (packed_ptr[word_idx] >> offset) & 1u; float sign = bit ? 1.0f : -1.0f; acc += static_cast(q_ptr[d]) * sign; } acc = simd_sum(acc); if (thread_index_in_simdgroup == 0) { auto idx = (b * kv_heads + h) * token_count + t; out[((b * kv_heads + h) * repeat_count + repeat_idx) * token_count + t] = acc * static_cast(norms[idx]) * static_cast(residual_norms[idx]) * scale[0]; } turboquant_qjl_scorerr)q_projr%residual_normssignsscaler(r+s r_qjl_score_kernelrasB   &FN 77   #KW  rcv\4'gR#Rp\PPR.ROR.VR7#)Na{ auto lane = thread_position_in_grid.x; auto repeat_idx = thread_position_in_grid.y; auto n = thread_position_in_grid.z; auto token_count = norms_shape[2]; auto kv_heads = norms_shape[1]; auto repeat_count = q_rot_shape[2]; if (repeat_idx >= repeat_count) { return; } auto b = n / (kv_heads * token_count); auto rem = n % (kv_heads * token_count); auto h = rem / token_count; auto t = rem % token_count; auto q_rot_ptr = q_rot + ((b * kv_heads + h) * repeat_count + repeat_idx) * Dim; auto q_proj_ptr = q_proj + ((b * kv_heads + h) * repeat_count + repeat_idx) * Dim; auto mse_ptr = mse_packed + ((b * kv_heads + h) * token_count + t) * MsePackedWidth; auto sign_ptr = signs + ((b * kv_heads + h) * token_count + t) * SignPackedWidth; float mse_acc = 0.0f; float qjl_acc = 0.0f; for (int d = lane; d < Dim; d += 32) { int bit_offset = d * MseBits; int word_idx = bit_offset / 32; int offset = bit_offset % 32; uint value = mse_ptr[word_idx] >> offset; int spill = offset + MseBits - 32; if (spill > 0) { value |= mse_ptr[word_idx + 1] << (MseBits - spill); } value &= ((1u << MseBits) - 1u); mse_acc += static_cast(q_rot_ptr[d]) * codebook[value]; int sign_word = d / 32; int sign_offset = d % 32; uint bit = (sign_ptr[sign_word] >> sign_offset) & 1u; float sign = bit ? 1.0f : -1.0f; qjl_acc += static_cast(q_proj_ptr[d]) * sign; } mse_acc = simd_sum(mse_acc); qjl_acc = simd_sum(qjl_acc); if (thread_index_in_simdgroup == 0) { auto idx = (b * kv_heads + h) * token_count + t; out[((b * kv_heads + h) * repeat_count + repeat_idx) * token_count + t] = static_cast(norms[idx]) * ( mse_acc + scale[0] * static_cast(residual_norms[idx]) * qjl_acc ); } turboquant_prod_scorerrr$r]r%r^ mse_packedr_r'r`r(r+s r_prod_score_kernelrf(sF   5Fl 77   $  W  rcv\4'gR#Rp\PPR.ROR.VR7#)NaV auto lane = thread_position_in_grid.x; auto dim_idx = thread_position_in_grid.y; auto n = thread_position_in_grid.z; if (dim_idx >= Dim) { return; } auto token_count = norms_shape[2]; auto kv_heads = norms_shape[1]; auto repeat_count = weights_shape[2]; auto b = n / (kv_heads * repeat_count); auto rem = n % (kv_heads * repeat_count); auto h = rem / repeat_count; auto repeat_idx = rem % repeat_count; auto weights_ptr = weights + ((b * kv_heads + h) * repeat_count + repeat_idx) * token_count; auto norms_ptr = norms + (b * kv_heads + h) * token_count; auto packed_ptr = packed + ((b * kv_heads + h) * token_count) * PackedWidth; float acc = 0.0f; for (int t = lane; t < token_count; t += 32) { auto token_ptr = packed_ptr + t * PackedWidth; int bit_offset = dim_idx * Bits; int word_idx = bit_offset / 32; int offset = bit_offset % 32; uint value = token_ptr[word_idx] >> offset; int spill = offset + Bits - 32; if (spill > 0) { value |= token_ptr[word_idx + 1] << (Bits - spill); } value &= ((1u << Bits) - 1u); acc += static_cast(weights_ptr[t]) * static_cast(norms_ptr[t]) * codebook[value]; } acc = simd_sum(acc); if (thread_index_in_simdgroup == 0) { out[((b * kv_heads + h) * repeat_count + repeat_idx) * Dim + dim_idx] = acc; } turboquant_mse_weighted_rotrrweightsr%r&r'r(r+s r_mse_weighted_rot_kernelrktsB   *FV 77   *>W  rcV^8dQhRR/#r.r)rs"rrrskkCkrc\4'dV^8:dR#.R!Op\V4F-pVPRV R24VPRV R24K/ V.R"O, p\V4F3pVPRV RV R 24VPR V R V R 24K5 V.R#O, p\V4FApVPRV R24VPR4VPRV RV R24KC VPR4\V4F3pVPRV RV R24VPRV RV R24K5 VRR., p\V4F0pVPRV R24VPRV RV R24K2 VR ., pRPV4p\P P RV 2.R$OR.VR 7#)%Nr2z float grp_z[num_entries] = {}; float qjl_acc_r6z grp_z)[value] += static_cast(q_rot_base[z * Dim + d]); qjl_acc_# += static_cast(q_proj_base[ * Dim + d]) * sign;r4z float mse_acc_z- for (int k = 0; k < num_entries; k++)z mse_acc_z += grp_z[k] * codebook[k];z float mse_sum_z = simd_sum(mse_acc_r; float qjl_sum_ = simd_sum(qjl_acc_- if (thread_index_in_simdgroup == 0) {5 out[((b * kv_heads + h) * repeat_count + ) * token_count + t] =z norm * (mse_sum_& + scale[0] * residual_norm * qjl_sum_r?turboquant_prod_score_repeat_rr)r@rAr2rBrC+ auto repeat_count = q_rot_shape[2];z7 constexpr uint mse_mask = (1u << MseBits) - 1u;z1 constexpr int num_entries = 1 << MseBits;r2. auto b = n / (kv_heads * token_count);0 auto rem = n % (kv_heads * token_count);# auto h = rem / token_count;# auto t = rem % token_count;r2zL auto q_rot_base = q_rot + ((b * kv_heads + h) * repeat_count) * Dim;zN auto q_proj_base = q_proj + ((b * kv_heads + h) * repeat_count) * Dim;z\ auto mse_ptr = mse_packed + ((b * kv_heads + h) * token_count + t) * MsePackedWidth;Y auto sign_ptr = signs + ((b * kv_heads + h) * token_count + t) * SignPackedWidth;r28 auto idx = (b * kv_heads + h) * token_count + t;4 float norm = static_cast(norms[idx]);F float residual_norm = static_cast(residual_norms[idx]);r2)r2. for (int d = lane; d < Dim; d += 32) {z) int bit_offset = d * MseBits;z+ int word_idx = bit_offset / 32;z) int offset = bit_offset % 32;z5 uint value = mse_ptr[word_idx] >> offset;z. int spill = offset + MseBits - 32; if (spill > 0) {zD value |= mse_ptr[word_idx + 1] << (MseBits - spill);r8z value &= mse_mask;r2# int sign_word = d / 32;% int sign_offset = d % 32;A uint bit = (sign_ptr[sign_word] >> sign_offset) & 1u;, float sign = bit ? 1.0f : -1.0f;r2)r4r2u6 // Dot product: group sums × codebook entriesrdrrOappendrPrr)r*r/rQrRr"s& r_prod_score_repeat_kernelrs4   !2  E4<  )!,ABC -aS9:! E(<  qc!J1#] [   "1#%HK_ ` ! E <  -aS9: DF +A3hqc9KLM!  LL <  -aS0DQCrJK -aS0DQCrJK!  7E<  CA3F\ ]   .qc1WXYWZZ\ ] ! EYYu F 77  ,\N ;  W  rcV^8dQhRR/#)r level_bitstuple[int, ...]r)rs"rrrsJJJrc\4'd\V4^8XdR#.R6Op\\V44FpVPRV^, 24K \\V44F9pVPRV^, 24VPRV^, 24K; .R7Op\ V^R7Fwr$^V,^, pVRV RV R V R 2R V R V R 2R V RV RV R 2RV RV R2RV RV R2RV RV RV RV R 2 RV RV RV R2RV R2RV R V RV R!V R"V R#2 R$R%V R&V R'2R(V R)V^, R*2R+V R,V R-V R.V R-V R/2 R., pK V.R8O, p\ P PR0R1PR2V44,VR3.R4PV4R57#)9rNangles_cos_sin_r2startz auto angle_ptr_z = angles_z7 + ((b * kv_heads + h) * token_count + t) * PackedWidthr=z int angle_idx_z = d >> z int bit_offset_z = angle_idx_z * z int word_idx_z = bit_offset_z / 32;z int offset_z % 32;z uint value_z = angle_ptr_z [word_idx_z ] >> offset_z int spill_z = offset_ +  - 32;z if (spill_z > 0) {z value_z |= angle_ptr_z + 1] << (z - spill_r;r8z value_z &= u;z bool use_sin_z = ((d >> z ) & 1) != 0;z coeff *= use_sin_z ? static_cast(sin_z[value_z]) : static_cast(cos_]);turboquant_polar_prod_score__c38"TFp\V4xK R#5iNstr.0bits& r +_polar_prod_score_kernel..bs6V:Cs3xx:rr?r)r$r%radii)r@z4 auto repeat_idx = thread_position_in_grid.y;rAr2rBrCrzz) if (repeat_idx >= repeat_count) { return;r4r2r{r|r}r~r2zT auto q_ptr = q_rot + ((b * kv_heads + h) * repeat_count + repeat_idx) * Dim;U auto radii_ptr = radii + ((b * kv_heads + h) * token_count + t) * BlockCount;r2z float acc = 0.0f;rz( int block_idx = d >> Levels;zC float coeff = static_cast(radii_ptr[block_idx]);r2) z8 acc += static_cast(q_ptr[d]) * coeff;r4r2z acc = simd_sum(acc);ruz< auto idx = (b * kv_heads + h) * token_count + t;zU out[((b * kv_heads + h) * repeat_count + repeat_idx) * token_count + t] =z5 acc * static_cast(norms[idx]);r4) rlenrOr enumeraterr)r*rP)rr levelrQbitsmasks& r_polar_prod_score_kernelrs{   ZA!5-Ks:'WUQYK01(s:'T%!)-.T%!)-.( E6!15 T Q )% 5'Axy~x@A B(xwa @)% eWCvQ O'wnUG6 J%eWN5' H%eWM% 5'Q]^c]dde f$UG:eWCvV L$UG8 4$UG>% 5'Q[\`[aajkpjqqs t  tD6 4'wj < P+E72LUGSZ[`Zaa}D~EELMRLSSV W   6&  E 77   +chh6V:6V.V VWyy  rc V^8dQhRRRR/#)r rrr/r0r)rs"rrrjs~~~PS~rcr\4'd\V4^8wgV^8:dR#Vwr#rE^V,^, p^V,^, p^V,^, p^V,^, p .RpOp \^4Fp V PRV ^, 24K V P .RqO4\^4F9p V PRV ^, 24V PRV ^, 24K; .RrOp \V4FWp V PRV R24V PRV R 24V PR V R 24V PR V R 24KY V R., p \V4Fp V PRV R24K V RRRV R2RRRRV R2RRV R2RRV R2RRR R!R., p \V4Fp V PR"V R#V R$V R%24K V R&R'RR(R)V R2R*R+R,R-V R2R.R/V R02RR1V R2R2R3R4R., p \V4Fp V PR5V R6V R7V R824K V R&R'RR9R:V R2R;RV R2R?R@V RA2RRBV R2RCRDR4R., p \V4Fp V PREV RFV RGV RH24K V R&R'RRIRJV R2RKRLRMRNV R2RORPV RQ2RRRV R2RSRTR4R., p \V4Fp V PRUV RVV RWV RX24K V .RsO, p \V4Fp V PRYV RZV R[24K V R&R., p \V4FKp V PR\V R]V R^24V PR_V R`V Ra24V PRbV RcV Ra24KM V RRd., p \V4F0p V PReV Rf24V PRgV RhV Ra24K2 V R&., p \ P PRiRjPRkV44,RlV 2,V Rm.RnPV 4Ro7#)tNrrrr2z! threadgroup float level1_z [Dim / 2];z! threadgroup float level2_z [Dim / 4];z! threadgroup float level3_z [Dim / 8];z! threadgroup float level4_z [BlockCount];ror6zG for (int pair_idx = lane; pair_idx < Dim / 2; pair_idx += 32) {z* int bit_offset_1 = pair_idx * r=z/ int word_idx_1 = bit_offset_1 / 32;z- int offset_1 = bit_offset_1 % 32;z> uint value_1 = angle1_ptr[word_idx_1] >> offset_1;z% int spill_1 = offset_1 + rz if (spill_1 > 0) {z: value_1 |= angle1_ptr[word_idx_1 + 1] << (z - spill_1);r8z value_1 &= rzA float cos_1_val = static_cast(cos_1[value_1]);zA float sin_1_val = static_cast(sin_1[value_1]);z# int d0 = pair_idx << 1;z int d1 = d0 + 1;z level1_z+[pair_idx] = static_cast(q_rot_base[z: * Dim + d0]) * cos_1_val + static_cast(q_rot_base[z * Dim + d1]) * sin_1_val;r48 threadgroup_barrier(mem_flags::mem_threadgroup);zG for (int pair_idx = lane; pair_idx < Dim / 4; pair_idx += 32) {z* int bit_offset_2 = pair_idx * z/ int word_idx_2 = bit_offset_2 / 32;z- int offset_2 = bit_offset_2 % 32;z> uint value_2 = angle2_ptr[word_idx_2] >> offset_2;z% int spill_2 = offset_2 + z if (spill_2 > 0) {z: value_2 |= angle2_ptr[word_idx_2 + 1] << (z - spill_2);z value_2 &= zA float cos_2_val = static_cast(cos_2[value_2]);zA float sin_2_val = static_cast(sin_2[value_2]);z& int child = pair_idx << 1;z level2_z[pair_idx] = level1_z[child] * cos_2_val + level1_z[child + 1] * sin_2_val;zG for (int pair_idx = lane; pair_idx < Dim / 8; pair_idx += 32) {z* int bit_offset_3 = pair_idx * z/ int word_idx_3 = bit_offset_3 / 32;z- int offset_3 = bit_offset_3 % 32;z> uint value_3 = angle3_ptr[word_idx_3] >> offset_3;z% int spill_3 = offset_3 + z if (spill_3 > 0) {z: value_3 |= angle3_ptr[word_idx_3 + 1] << (z - spill_3);z value_3 &= zA float cos_3_val = static_cast(cos_3[value_3]);zA float sin_3_val = static_cast(sin_3[value_3]);z level3_z[pair_idx] = level2_z[child] * cos_3_val + level2_z[child + 1] * sin_3_val;zJ for (int pair_idx = lane; pair_idx < BlockCount; pair_idx += 32) {z* int bit_offset_4 = pair_idx * z/ int word_idx_4 = bit_offset_4 / 32;z- int offset_4 = bit_offset_4 % 32;z> uint value_4 = angle4_ptr[word_idx_4] >> offset_4;z% int spill_4 = offset_4 + z if (spill_4 > 0) {z: value_4 |= angle4_ptr[word_idx_4 + 1] << (z - spill_4);z value_4 &= zA float cos_4_val = static_cast(cos_4[value_4]);zA float sin_4_val = static_cast(sin_4[value_4]);z level4_z[pair_idx] = level3_z[child] * cos_4_val + level3_z[child + 1] * sin_4_val;rprqrrz float polar_acc_zD = lane < BlockCount ? static_cast(radii_ptr[lane]) * level4_z[lane] : 0.0f;z float polar_sum_z = simd_sum(polar_acc_r;rsrtruz4 out[((b * kv_heads + h) * RepeatCount + rwz" norm * (polar_sum_rxturboquant_polar_turbo_score_rc38"TFp\V4xK R#5irrrs& rr3_polar_turbo_score_repeat_kernel..#s2z3s88zr_repeat_rr?r)r$r]r%r)r^r_r`)r@rAr2rBrCr2r{r|r}r~r2zK auto q_rot_base = q_rot + ((b * kv_heads + h) * RepeatCount) * Dim;zM auto q_proj_base = q_proj + ((b * kv_heads + h) * RepeatCount) * Dim;rz[ auto angle1_ptr = angles_1 + ((b * kv_heads + h) * token_count + t) * PackedWidth1;z[ auto angle2_ptr = angles_2 + ((b * kv_heads + h) * token_count + t) * PackedWidth2;z[ auto angle3_ptr = angles_3 + ((b * kv_heads + h) * token_count + t) * PackedWidth3;z[ auto angle4_ptr = angles_4 + ((b * kv_heads + h) * token_count + t) * PackedWidth4;rr2rrrr2) r4rr2rrrrrr2) rrrOrextendrr)r*rP)rr/bits1bits2bits3bits4mask1mask2mask3mask4r rrQrRs&& r _polar_turbo_score_repeat_kernelris   ZA!59J!+E% %Z1 E %Z1 E %Z1 E %Z1 E7KqWUQYK01;<qT%!)-.T%!)-. E4<  8:FG 8:FG 8:FG 8=IJ !  bTME <  -aS9:!  Q 4UG1=97H /wf=( DUG<X !%+KK-& !E$<  !!$OPQsSMNOMPPj k ! B Q 4UG1=97H /wf=( DUG<X !%+KK0 #E&<  !!$8;XYZX[[s t ! B Q 4UG1=97H /wf=( DUG<X !%+KK0 #E&<  !!$8;XYZX[[s t ! B T 4UG1=97H /wf=( DUG<X !%+KK0 #E&<  !!$8;XYZX[[s t !  E<  "1#%HK_ ` !  E<  &qc)mnompp~    /s2H2NO -aS0DQCrJK !   7E<  B1#E[ \   03YZ[Y\\^ _ ! E 77   , ((2z2 23 \N #$ Wyy  rcV^8dQhRR/#r.r)rs"rrr,sGG#GrcP\4'dV^8:dR#.ROp\V4FpVPRV R24K V.RO, p\V4FpVPRV RV R24K VRR., p\V4FpVPR V R V R 24K VRR ., p\V4FpVPR V RV R24K VR., pRPV4p\P P RV 2.ROR.VR7#)rnNr2r4 float acc_r6 acc_z$ += static_cast(weights_base[z# * token_count + t]) * norm * code; float acc_sum_ = simd_sum(acc_r;rurv) * Dim + dim_idx] = acc_sum_r=r?#turboquant_mse_weighted_rot_repeat_rr)r@1 auto dim_idx = thread_position_in_grid.y;rAr2 if (dim_idx >= Dim) {rr4r2rBrCz- auto repeat_count = weights_shape[2]; auto b = n / kv_heads; auto h = n % kv_heads;r2zX auto weights_base = weights + ((b * kv_heads + h) * repeat_count) * token_count;B auto norms_ptr = norms + (b * kv_heads + h) * token_count;T auto packed_ptr = packed + ((b * kv_heads + h) * token_count) * PackedWidth;r2( int bit_offset = dim_idx * Bits;' int word_idx = bit_offset / 32;% int offset = bit_offset % 32;r2 r2z6 for (int t = lane; t < token_count; t += 32) {z: auto token_ptr = packed_ptr + t * PackedWidth;z7 uint value = token_ptr[word_idx] >> offset;z+ int spill = offset + Bits - 32;rzC value |= token_ptr[word_idx + 1] << (Bits - spill);r8z) value &= ((1u << Bits) - 1u);z) float code = codebook[value];z: float norm = static_cast(norms_ptr[t]);rirrs& r_mse_weighted_rot_repeat_kernelr+se   !2 E0<  )!H56!  E<  qc!EaSHk l !  E<  -aS0@2FG!  7E<  CA3Fcdecffg h ! EYYu F 77  2<. A>W  rcV^8dQhRR/#r.r)rs"rrrwsTTTrc j\4'dV^8:dR#.R!Op\V4FEpVPRV RV R24VPRV R24VPR V R24KG V.R"O, p\V4FNpVPR V R V R V R 24VPRV RV R24VPRV RV R24KP VRR., p\V4F3pVPRV RV R 24VPRV RV R 24K5 VRR., p\V4F0pVPRV R24VPRV RV R24K2 VR., pRPV4p\P P RV 2.R#OR.VR 7#)$zySingle-pass fused softmax + weighted sum kernel. Takes precomputed max_scores to avoid a separate token-dimension pass. Nr2r4 float max_score_ = static_cast(max_base[rrr6z float denom_ float weight_& = exp(static_cast(scores_base[! * token_count + t]) - max_score_r;r += weight_ * norm * code;z denom_r=rrz float denom_sum_z = simd_sum(denom_rurvz) * Dim + dim_idx] =z acc_sum_z / max(denom_sum_z , 1e-6f);r?*turboquant_mse_scores_weighted_rot_repeat_rrr@rrAr2rrr4r2rBrCz, auto repeat_count = scores_shape[2];rrr2zV auto scores_base = scores + ((b * kv_heads + h) * repeat_count) * token_count;rrzG auto max_base = max_scores + (b * kv_heads + h) * repeat_count;r2rrrr2rscoresr%r&r' max_scoresrrs& r&_mse_scores_weighted_rot_repeat_kernelrvs   !2 E2<  &qc)H3 O   )!H56 +A3h78 !   E<  's*PQRPSStuvtwwy z   's+aSHI )!Ks!<= !   E<  -aS0@2FG /s2DQCrJK!  7E<  CA3FZ [   /s2CA3iPQ !  EYYu F 77  9, HKW  rcV^8dQhRR/#r.r)rs"rrrsPPSPrc \4'dV^8:dR#.ROp\V4F0pVPRV RV R24VPRV R24K2 V.RO, p\V4F6pVPR V R V R V R 24VPR V RV R24K8 VRR., p\V4FpVPRV RV R 24K VRR., p\V4FpVPRV RV R24K VR., pRPV4p\P P RV 2.ROR.VR7#)zSingle-pass kernel for unnormalized weighted sum (used in chunked attention). Takes precomputed max_scores to avoid a separate token-dimension pass. Nr2r4rrrrr6rrrr;rrrrrrurvrr=r?.turboquant_mse_scores_weighted_rot_sum_repeat_rrrrrrrs& r*_mse_scores_weighted_rot_sum_repeat_kernelrs   !2 E2<  &qc)H3 O   )!H56 !   E<  's*PQRPSStuvtwwy z   's+aSHI !   E<  -aS0@2FG!  7E<  CA3Fcdecffg h ! EYYu F 77  =l^ LKW  rc ,V^8dQhRRRRRRRRRR /#) r r$mx.arraystateTurboQuantMSEStaterr0r'r Optional[mx.array]r)rs"rrr"s:8*8* 8* 8* 8* 8*  8*rc V^8:gD\4'd4VP^8wg#VPP^,^8XdR#VPwrErgVPP^,pV^,^ ,p \ V4p V e^@p W,^, V ,p V !WPVP V.RV3RV 3RV3RVP PR ,3RV3RV 3.^ ^WE,V ,3R WEWh3.\ P.R7^,p \ P!V ^R 7#\4pVfR#V!WPVP V.RV3RV3RVP PR ,3.^ WdV,V,3R WEWh3.\ P.R7^,p \ P!V ^R 7#) rNDim DimsPerLaneBits PackedWidth RepeatCount TokTileSizeinputstemplategrid threadgroup output_shapes output_dtypesaxis rnrn) rndimr%shaperSindicesrfloat32 expand_dimsr,)r$rrr'BHRDT dims_per_lane tiled_kernel tok_tile_size num_tok_tilesrkernels&&&& r_metal_mse_scorer"s  !! ::? ;;  Q 1 $JA! !AVNM+1-L *Q.=@ ;; x@  . 3 3B 78" . a./"!<.::,   ~~f1--  F ~ {{EMM8< AJ TN EMM//3 4 !UQY a|nzzl   F >>&q ))rc(V^8dQhRRRRRRRR/#)r r]rrTurboQuantProdStater`r rr)rs"rrr]s0** * * * *rc *\4'd4VP^8wg#VPP^,^8XdR#\ 4pVfR#VPwrErgVPP^,pV!VVPVP VP V.RV3RVP PR,3.^ WdV,V,3RWEWh3.\P.R7^,p \P!V ^R7#)rNrrrrrr) rrr%rrar^ qjl_signsrrr) r]rr`r rrrrrrs &&& r_metal_qjl_scorer]s   !1U[[5F5Fq5IQ5N  F ~JA! !A   KK  OO   AJ EOO11"5 6 !UQY a|nzzl ! F" >>&q ))rc4V^8dQhRRRRRRRRRRR RR R /#) r r$rr]rrmse_bitsr0r'r`r rr)rs"rrrsWI*I* I* I* I* I*  I*  I*I*rc jV^8:gU\4'dEVP^8wg4VP^8wg#VPP^,^8XdR#VPwrgrVPP^,p V^8d\ V4p V eV !VVVPVP VP VPVV.RV 3RV3RV3RVP PR ,3RVPPR ,3.^ ^Wg,V ,3R WgW3.\P.R7^,p \P!V ^R7#\4p V fR#V !VVVPVP VP VPVV.RV 3RV3RVP PR ,3RVPPR ,3.^ WV,V ,3R WgW3.\P.R7^,p \P!V ^R7#) rNrrMseBitsMsePackedWidthSignPackedWidthrrrr) rrr%rrr^ mse_indicesrrrrrf) r$r]rrr'r`rrrrrr rs &&&&&& r_metal_prod_scorers A !! ::? ;;!  ;;  Q 1 $JA! !A1u*1-  KK((%%OO AJ"A&)%u'8'8'>'>r'BC&(=(=b(AB !QUQY'& !a|n!zzl+,-F.>>&q1 1  !F ~    KK     OO    AJ  ! u0066r: ;  5 5b 9 :  !UQY a|nzzl)*+ F, >>&q ))rc 0V^8dQhRRRRRRRRR RR R /#) r r$rrTurboQuantPolarProdStaterr cos_tablestuple[mx.array, ...] sin_tablesr rr)rs"rrrsD+*+* +* #+* +*% +* % +*  +*rc \4'dDVP^8wg3VPP^,^8Xg\ V4^8XdR#\ V4pVfR#VPwrgrVPP^,p \ V4p WPVP P.p V PRVP P44\W44FwrV PW.4K RV 3RV 3RVP PPR ,3.p\VP P^R7F-wppVPRV 2VPR ,34K/ V!V V^ WV,V ,3R WgW3.\P.R7^,p\P!V^R 7#) rNc3$"TFqxK R#5irrrrs& rr*_metal_polar_prod_score..E%DE%%DrLevels BlockCountrrrrrr)rrr%rrr polar_staterr level_indicesziprrrrr)r$rrrr r rrrrrlevelsr cos_table sin_tabler level_idxrrs&&&&& r_metal_polar_prod_scorer0s    ::? ;;  Q 1 $ z?a  %j 1F ~JA! !A _F[[%"3"3"9"9 :F MMEU%6%6%D%DEE #J ;  y,-!<  6 u((..44R89H &e&7&7&E&EQO 5;yk2EKKODEP!UQY a|nzzl  F >>&q ))rc8V^8dQhRRRRRRRRRR R R R RR R /#)r r$rr]rrrrrrr r`r rr)rs"rrrsX8*8* 8* 8* $8* 8* % 8* % 8* 8*8*rc \4'dpVP^8wg_VP^8wgNVPVP8wg3VPP^,^8Xg\ V4^8XdR#VPwrxrVPP^,p \ V4p \ W94p V fR#WVPVP P.pVPRVP P44VPVPVPV.4\WE4FwppVPVV.4K RV 3RV 3RV 3RVP PPR ,3RVPPR ,3.p\VP P^R7F-wppVPR V 2VPR ,34K/ V !VV^ ^Wx,V ,3R WxW3.\P .R 7^,p\P"!V^R 7#)rNc3$"TFqxK R#5irrr#s& rr+_metal_polar_turbo_score..r%r&rr'rr(rrrrrrr)rrrr%rrr)rrr*r^rr+rrrrr)r$r]rrrr r`rrrrrr,r rr-r.rr/rrs&&&&&&& r_metal_polar_turbo_scorer5s     ::? ;;!  ;;&,, & ;;  Q 1 $ z?a JA! !A _F -j >&q ))rc(V^8dQhRRRRRRRR/#)r weighted_rotrrotationr_rr r)rs"rrr4s,---&.-7I- -rcLVe \W4#\P!W4#)aUndo the codec rotation on a Metal value-kernel result. The L=1 value kernels return the weighted value sum in the codec's *rotated* space. Map it back to model space with the same inverse the codec uses: an explicit Randomized Hadamard inverse when the codec runs with RHT (``signs`` provided), otherwise the dense-rotation matmul. Mirrors ``_TurboQuantMSECodec._rotate_inverse`` so the kernel fast path stays valid under RHT instead of being skipped for the einsum fallback. ) _rht_inversermatmul)r7r8r_s&&&r_value_rotate_inverser<4s$ L00 99\ ,,rc4V^8dQhRRRRRRRRRRR R R R /#) r rjrrrrr0r'r8r_rr r)rs"rrrEsWG*G* G* G* G* G*  G*  G*G*rc ^V^8:g\\4'dLVP^8wg;VPR,^8wg#VPP^,^8XdR#VP VP^,VP^,VP^,VPR ,4pVPwrxrVP^,p V ^8d\ V 4p V eV !VVPVP V.RV 3RV 3RV3RVP PR ,3.^ WV,3R WxW3.\P.R7^,p \WV4p\P!V^R7#\4p V fR#V !VVPVP V.RV 3RV3RVP PR ,3.^ WV,V ,3R WxW3.\P.R7^,p \WV4p\P!V^R7#) rNrrrrrrrr) rrrr%reshaperrrrr<rrk)rjrrr'r8r_ weights_2drrrrrr r7outputs&&&&&& r_metal_mse_weighted_sumrCEs  !! <<1  ==  ! ;;  Q 1 $ a a a b J !!JA!qA1u03  !KKMM AJ"A&TN"EMM$7$7$;<  !U^& !a|n!zzl!"#L$+<5IF>>&q1 1 % 'F ~  KK MM   AJ TN EMM//3 4 !UQY a|nzzl ! L"#<5 AF >>&q ))rc4V^8dQhRRRRRRRRRRR R R R /#) r rrrrrr0r'r8r_rr r)rs"rrrsN6*6* 6* 6* 6* 6*  6*  6*6*rc DV^8:g\\4'dLVP^8wg;VPR,^8wg#VPP^,^8XdR#VP VP^,VP^,VP^,VPR,4pVPwrxrV ^8:dR#\ V 4p V fR#\ P!VRR7p VP^,p V !VVPVPVV .RV 3RV3RVPPR,3.^ WV,3R WxW3.\ P.R7^,p\WV4p\ P!V^R7#) rNrrrrrr?rr) rrrr%r@rrmaxrrr<r)rrrr'r8r_ scores_2drrrrr rrr7rBs&&&&&& r#_metal_mse_weighted_sum_from_scoresrHsn  !! ;;!  << q ;;  Q 1 $ Q Q Q R I JA!Av 3A 6F ~ +JqA  KK MM    AJ TN EMM//3 4 !U^a|nzzl!"# L$#<5 AF >>&q ))rc8V^8dQhRRRRRRRRRRR RR R R R /#) r rrrrrr0r'r8rr_rr r)rs"rrrsX5*5* 5* 5* 5* 5*  5*  5* 5*5*rc V^8:g\\4'dLVP^8wg;VPR,^8wg#VPP^,^8XdR#VP VP^,VP^,VP^,VPR,4pVPwrrV ^8:dR#\ V 4p V fR#VP^,p V !VVPVP VV.RV 3RV3RVP PR,3.^ WV ,3R WW3.\P.R7^,p\WV4p\P!V^R7#) rNrrrrrr?rr) rrrr%r@rrrrr<r)rrrr'r8rr_rGrrrrr rr7rBs&&&&&&& r'_metal_mse_weighted_sum_sum_from_scoresrKs^  !! ;;!  << q ;;  Q 1 $ Q Q Q R I JA!Av 7 :F ~ qA  KK MM    AJ TN EMM//3 4 !U^a|nzzl!"# L$#<5 AF >>&q ))rcV^8dQhRR/#r rr0r)rs"rrrs//#/rcjaa\S^, ^4o\PRVV3Rll4pV#)rncLV^8dQhRRRRRRRRRRRRRRR RR RR RR RR RRR/ #)r grouped_queriesr key_normskey_mse_indiceskey_residual_norms key_qjl_signs value_norms value_indiceskey_query_transform_t key_codebook key_scalevalue_codebookvalue_rotationr r)rs"rr5_compiled_integer_decode_kernel..__annotate__s) ) !) ) ") % )  )  )  )  () ) ) !) !)  ) rc J<\P!W4p VPR,p V RRV 13,pV RV R13,p\VP VP^,VP^,VP^,VPR,4VP VP^,VP^,VP^,VPR,4\ VVVV4SVV 4p\ V\WV4SV V 4#)rn.Nr)rr;rrr@rrHr)rPrQrRrSrTrUrVrWrXrYrZr[query_transformeddimr$r]rrrs&&&&&&&&&&&& r_decode0_compiled_integer_decode_kernel.._decodesIIoM##B'!#tt),"39-" MM A A A B  NN Qa&,,q/6<R?R@RAV RB2R8., p\V4FQpVRCV RDV R)2REV RFV RGV R;2RHV RIV RJV R;2RKV RLV RJV R;2RMV RNV ROV RPV R)2 RQV RRV ROV RPV RS2 ., pKS VRTR., pVRU., p\V4F'pVRVRWV RXV R)2RTRYRZV R)2R[V R)2R\V R)2RT., pK) \PP R]V R^V 2.RbO.RcOR_P V4R`7#)dz 32); auto qr_! = q_rot + (base * RepeatCount + ) * Dim; auto qp_" = q_proj + (base * RepeatCount +  float lmax_ = -INFINITY; float lsum_r6 float lacc_r52 for (int d = lane; d < Dim; d += 32) { int bo = d * r=9 uint idx = (mse_t[bo >> 5] >> (bo & 31));! if (((bo & 31) + () > 32) idx |= mse_t[(bo >> 5) + 1] << ( - ((bo & 31) +  - 32)); idx &= r/ float code = key_codebook[idx];< uint sb = (sign_t[d >> 5] >> (d & 31)) & 1u;r7 += kn * static_cast(qr_0[d]) * code + ksr * (sb ? static_cast(qp_[d]) : -static_cast(qp_[d]));r8r9r:r; float v_code = 0.0f; if (v_valid) {z6 auto vt = v_packed + t * VPackedWidth;z0 uint vv = (vt[v_word] >> v_off);zY if (v_spills) vv |= vt[v_word + 1] << (ValBits - (v_off + ValBits - 32));z+ v_code = val_codebook[vv & z$u] * static_cast(v_norms[t]); float om_ = lmax_ lmax_ = max(lmax_, s_ float rs_ = exp(om_ - lmax_ float w_ = exp(s_ lsum_ = lsum_ * rs_ + w_ lacc_ = lacc_ * v_code;r4z int out_stride = Dim; if (v_valid) {E out_acc[((bh * num_tok_tiles + tok_tile) * RepeatCount + !) * out_stride + val_dim] = lacc_ if (val_dim == 0) {J int sm_base = (bh * num_tok_tiles + tok_tile) * RepeatCount + $ out_sum[sm_base] = lsum_$ out_max[sm_base] = lmax_ turboquant_fused_integer_decode__rr?r)r2rL5 auto mse_t = k_mse + t * KMsePackedWidth;9 auto sign_t = k_signs + t * KSignPackedWidth;6 float kn = static_cast(k_norms[t]);I float ksr = kn * key_scale[0] * static_cast(k_res[t]); r$r]rQkey_mse key_res_norms key_signs val_norms val_packedrXrY val_codebookout_accout_sumout_maxrrFrOrr)r*rP)rr/rfrmse_maskval_maskrQrRs&&& r_fused_integer_decode_kernelr3s   !1+q0|c$(A6FHX "HT QH 8 <  6   6   9  6  V  /  5  *  *  2    8! " L# $ % & A' ( O) * C+ , T- . A/ 0 R1 2 3 4 05 6 07 8 *9 : ); < := > ? ED<  qc!B1#X Nqc!CA3h O  ! <  !!M 2!!H -!!H -  ! E<  )!H566! = 'z3 C +H:5]^f]ggwyAxBBJ K !(2. 9 FE<  !!$CA3FvwxvyzXYZX[[a b  ! E<  (?1#R@AA!  *$@: c 5hZ?cd  E<  #A3hqc 3s,qcaS ;#A3j8A3b A"1#Yqc!B ?s(1#VA3eA3a @s(1#VA3eA3j I   ! k2 E - ..E <  %STUSVVwxywzz{ |*XYZX[[\ ]21#Q 721#Q 7  ! 77  /vR ~ F  7yy!  rc(V^8dQhRRRRRRRR/#)r rfr0r/ num_queriesr r)rs"rrrs2AAA%(A7:AKNArc"\4'dV^8gV^8dR#^V,^, p\PPRV RV RV 2.ROR.RV RV RV R V R V R V R V R V RV RV RV RV RV RV R2R7#)zMulti-query score kernel: unpack key data ONCE per token, loop over L queries. Avoids R*L repeat explosion that causes register spill.Nmq_prod_score_r_lra auto lane = thread_position_in_grid.x; auto ri = thread_position_in_grid.y; auto n = thread_position_in_grid.z; auto tc = key_norms_shape[2]; auto kh = key_norms_shape[1]; auto b = n / (kh * tc); auto rem = n % (kh * tc); auto h = rem / tc; auto t = rem % tc; if (ri >= al) return; auto mt = key_mse + ((b*kh+h)*tc+t) * KMsePackedWidth; auto st = key_signs + ((b*kh+h)*tc+t) * KSignPackedWidth; float kn = static_cast(key_norms[(b*kh+h)*tc+t]); float ksr = kn * key_scale[0] * static_cast(key_res_norms[(b*kh+h)*tc+t]); // Unpack key ONCE float kc[z], ksf[zY]; for (int i=0, d=lane; d < Dim; i++, d+=32) { int bo = d * zR; uint idx = (mt[bo>>5] >> (bo&31)); if (((bo&31)+z!) > 32) idx |= mt[(bo>>5)+1] << (z - ((bo&31)+z-32)); idx &= zu; kc[i] = key_codebook[idx]; uint sb = (st[d>>5] >> (d&31)) & 1u; ksf[i] = sb ? 1.0f : -1.0f; } // Loop over L queries, reusing unpacked key auto bq = (b*kh+h) * z' + ri; for (int l = 0; l < z; l++) { float ps = 0.0f; for (int i=0, d=lane; d < Dim; i++, d+=32) { ps += kn * static_cast(q_rot[(bq*zR+l)*Dim+d]) * kc[i] + ksr * static_cast(q_proj[(bq*z+l)*Dim+d]) * ksf[i]; } float s = simd_sum(ps); if (lane == 0) out[((b*kh+h)*z+ri)*z+*tc + l*tc + t] = s; } r)r$r]rQrrrrXrYr()rfr/rr rs&&&& r_multi_query_prod_score_kernelrs   !1[1_\!Q&H 77  l^2l^2k] K  W $n%$_GM?;*^,*^+L\NZfgsftu z""#/0!!, .=>IMJ??JmL#$0.k]C O)  7 7rc$V^8dQhRRRRRR/#)r rr0r/r r)rs"rrr s$== = =14=rc\4'dV^8dR#^V,^, pRRRRRRRR R R R R RR RV R2RRRV R2R .p\V4FpVRV R2., pK VR RRRRV RV R2RV R2R ., p\V4FpVRV R V R!2., pK VR"R ., p\V4FpVR#V R$V R2., pK \PP R%V R&V 2.R*OR'.R(P V4R)7#)+zSingle-tile value weighted sum with precomputed softmax weights. TG=Dim: one thread per value dim, no exp() calls in inner loop. 2x faster than online-softmax variant.Nz- auto dim = thread_position_in_grid.x;rArBrCrDrErFrGrHr2z; auto wt = weights + bh * RepeatCount * token_count;rJrIz int bo = dim * r=z int v_word = bo / 32;z int v_shift = bo % 32;z" bool v_spill = (bo % 32 + ) > 32;rr6rLrMz. uint vv = (pt[v_word] >> v_shift);z0 if (v_spill) vv |= pt[v_word+1] << (z - (v_shift+-32));z& float val = codebook[vv & zu] * static_cast(nm[t]);rz += wt[z*token_count+t] * val;r4zD if (dim < Dim) out[((bh*num_tok_tiles+tok_tile)*RepeatCount+z)*Dim+dim] = acc_turboquant_single_tile_value_rrr?rrirN)rr/r rrQrRs&&& r&_single_tile_value_weighted_sum_kernelr s   !1T QH 8541U.47K E5D !$q) ' ( ,TF': ' E,<  &qc233!  958 :4& TFRXY 0 :YZ E<  $QCwqc1GHII! k2 E <  RSTRUUfghfiij k  ! 77  ,TF"\N C>Wyy  rc(V^8dQhRRRRRRRR/#)r rr0r/r rfr)rs"rrrKs2UU U U14UDGUrc\4'dV^8dR#V^8dTM\V^, ^4p^V,^, p^V,^, p.RNRNRNRNRNRNRNR NR NR NR NR NRNRNRNRNRNRNRNRNRNRNRNRNRNRV R2NRNRNRV R2NR NRNp\V4FpVR!V R"V R#2R$V R%V R#2., pK \V4FpVR&V R'V R(2R)V R*2., pK V.RbO, p\V4FpVR+V R(2., pK VR,R-V R2R.R/V R0V R1V R22R3V R42R5R6., p\V4FpVR7V R8V R9V R:V R;2 ., pK VR<., p\V4FpVR=V R>V R?2., pK V.RcO, p\V4FLpVR@V R2RAV RBV RCV R?2RDV R?2REV RFV R?2RGV RHV RI2RJRKV RL2RMRNROV RPV RQ2RRV RSV RT2RURVR<., pKN VR R., p\V4F'pVRRWV RXV RY2R RZR[V R2R\V R2R]V R2R ., pK) \PP R^V R_V 2.RdO.ReOR`P V4Ra7#)fuSingle-tile fused kernel — each lane handles multiple value dims. Zero key read redundancy: keys are read once per token, not once per val_tile. Faster than multi-tile at long contexts (256k+) where memory bandwidth dominates. Nr@rArjrkrDrErFrlrmrnrGrHr2rorprqrrrsrtzF // Precompute value bit offsets for all dims this lane handles6 int v_words[DimsPerLane], v_offs[DimsPerLane];: bool v_spills[DimsPerLane], v_valids[DimsPerLane];D for (int i = 0, vd = lane; i < DimsPerLane; i++, vd += 32) {# v_valids[i] = vd < Dim; int vbo = vd * r=" v_words[i] = vbo / 32;! v_offs[i] = vbo % 32;& v_spills[i] = (vbo % 32 + rr4rwrxryrzr{r| = -INFINITY, lsum_r6r[DimsPerLane] = {};r5rrrrrrrrrrrr7rrrrr8r9r:r; { float om = lmax_ lmax_rr' float rs = exp(om - lmax_ float w = exp(s_r lsum_r * rs + w;5 for (int i = 0; i < DimsPerLane; i++) { lacc_ [i] *= rs;$ if (v_valids[i]) {> uint vv = (vt[v_words[i]] >> v_offs[i]);B if (v_spills[i]) vv |= vt[v_words[i]+1] << ( - (v_offs[i]+r lacc_[i] += w * val_codebook[vv & u] * vnorm; } }zK if (vd < Dim) out_acc[((bh*num_tok_tiles+tok_tile)*RepeatCount+z)*Dim+vd] = lacc_[i];z if (lane == 0) {zB int sm_base = (bh*num_tok_tiles+tok_tile)*RepeatCount+rr%turboquant_fused_integer_single_tile_rr?r)r2rLrrrrr2r2z2 auto vt = v_packed + t * VPackedWidth;z9 float vnorm = static_cast(v_norms[t]);rrr) rr/r rfrrrrQrRs &&&& r(_fused_integer_decode_single_tile_kernelrJs   !1+q0|c$(A6FHX "HT QH 8 5  9  6  V  /  5  *  *  2  8  L    A  O  C! " T# $ A% & R' ( ) * Q+ , A- . E/ 0 O1 2 .3 4 &dV1-5 6 -7 8 ,9 : 1g>; < = > ? ED<  qc!B1#X Nqc!CA3h O  ! <  !!$7s( C!!$9 :  !  E<  )!H566! = 'z3 C +H:5]^f]ggwyAxBBJ K !(2. 9 FE<  !!$CA3FvwxvyzXYZX[[a b  ! o E <  (?1#R@AA! E <  -aS 2!!L4s" =5aS ;,QCxs" =!!HQCz :D%aS 33LPQUPVVdeidjjp q)!,I(S^ _"  !$ k2 E<  SYZ[Y\\mnomppt u'PQRPSST U21#Q 721#Q 7  ! 77  4TF"\N K  7yy!  rc ,V^8dQhRRRRRRRRRR/#)r rr0r/r rfuse_rhtrr)rs"rrrsANN NNN N  Nrc\4'dV^8dR#V^8dTM\V^, ^4p^V,^, p^V,^, p^V,pRRRRRRRR R R R R RRR RRRRRRV R2RRRV R2RR .p \V4Fp V RV RV R2RV R V R2., p K \V4Fp V R!V R"V R#2R$V R%2., p K V .RO, p \V4Fp V R&V R#2., p K V R'R(V R2R)R*V R+V R,V R-2R.V R/2R0R1., p \V4Fp V R2V R3V R4V R5V R62 ., p K V R7., p \V4Fp V R8V R9V R:2., p K V .RO, p \V4FLp V R;V R2RV R:2R?V R:2R@V RAV R:2RBV RCV RD2RERFV RG2RHRIRJV RKV RL2RMV RNV RO2RPRQR7., p KN V RR ., p \V4Fp V RRV RSV R2RTRUV RVV R2., p K V'djV .RO, p \V4FNp V RWV 2RXRYV RZ2., p V^8gV .RO, p V \R[R\R]R^4, p V R_R`V Ra2R ., p KP MjV R RbV Rc2., p \V4Fp V RdReV RfV Rg2., p K V .RO, p \V4Fp V RRhRiRjRkRlV Rm2RnV Ro2R7R. , p K V'dRpMR p V'dR\MRqp \P P RrV RsV V 2RtRuRvRwRxRyRzR{R|R}R~V . R.RPV 4R7#)zFully fused decode: score + online softmax + value + normalize + rotation in a single Metal dispatch. Processes ALL tokens (no tiling), outputs in original space. Reduces dispatch count from 7 to 1.Nr@, auto bh = thread_position_in_grid.z;rjrkrlrmrnr2rorprqrrrsrtz' // Precompute value bit offsetsrrrrrr=rrrrr4rwrxryrzr{r|rr6rrr5rrrrrrrrrrrr7rrrrr8r9r:r;rrrrrrrrrrrrrrrrrrrrrrrz float inv_sum_z = 1.0f / lsum_- for (int i = 0; i < DimsPerLane; i++)rz[i] *= inv_sum_z" // RHT inverse for repeat zH for (int i = 0, vd = lane; i < DimsPerLanePadded; vd += 32, i++)z> shared_rot[vd < Dim ? vd : 0] = (vd < Dim) ? lacc_z [i] : 0.0f; shared_rotsign_vecwht_tmpDimsPerLanePaddedzB for (int i = 0, vd = lane; i < DimsPerLane; vd += 32, i++)z2 if (vd < Dim) out[(bh * RepeatCount + z) * Dim + vd] = shared_rot[vd];z% threadgroup float shared_rot[z * Dim];zB for (int i = 0, vd = lane; i < DimsPerLane; i++, vd += 32)z% if (vd < Dim) shared_rot[z * Dim + vd] = lacc_rz if (vd < Dim) {z% float rot_val = 0.0f;z5 auto rot_row = rotation_t + vd * Dim;2 for (int j = 0; j < (int)Dim; j++)z* rot_val += shared_rot[z * Dim + j] * rot_row[j];r<z) * Dim + vd] = rot_val;_rht rotation_tturboquant_fully_fused_rr$r]rQrrrrrrXrYrrr?r)r2z4 for (int t = 0; t < (int)token_count; t++) {rrrrr2r)r2z0 threadgroup float shared_rot[DimPadded];z) float wht_tmp[DimsPerLanePadded];)A if (lane < (DimPadded - Dim) && (Dim + lane) < DimPadded)z* shared_rot[Dim + lane] = 0.0f;r)rr2z9 // Apply dense rotation: output in original space)rrFrO_metal_butterfly_wht_inverserr)r*rP) rr/r rfrrrrnum_mse_entriesrQrR name_suffixrotation_inputs &&&&& r_fully_fused_decode_kernelrs   !1+q0|c$(A6FHX "HT QH8mO 9685))1 @NBS@Q 1@DN - %dV1- , + 0g> 5 E:<  qc!B1#X Nqc!CA3h O  ! <  !!$7s( C!!$9 :  ! E<  )!H566! = 'z3 C +H:5]^f]ggwyAxBBJ K !(2. 9 FE<  !!$CA3FvwxvyzXYZX[[a b  ! o E <  (?1#R@AA! E <  -aS 2!!L4s" =5aS ;,QCxs" =!!HQCz :D%aS 33LPQUPVVdeidjjp q)!,I(S^ _"  !$ k2 E<  $QCqc ;;s/!A 6  !    |$A 4QC8ZPQRPSS^_ E a4 1j)5H E TDQCGfg E!%.  3L> J  |$A T7s:NqcQUV E%    |$A W.7GDA4 4rc|^V,^, p.pV'g\V4FpW`,pV^,pV^,p W,^8:d#V ^8Xd V RV RV R2p MZV RV RV RV R2p MK^V , p ^W , ,^, p V RV RV R^V ,^, RV^, RV R V R 2p VPV 4K V#\V4FYpV'dR V R W`, R 2MTp R V R2pR V R2pV RV R V RV RV RV RV RV RV RV R2p VPV 4K[ V#)uGenerate fully-unrolled byte extraction for n_elems packed at `bits` bits. When *bit_off_var* is empty the extraction assumes element 0 starts at bit 0 of the byte pointer (fast path — compile-time constant offsets, works when ``n_elems * bits % 8 == 0``). When *bit_off_var* names a Metal ``int`` variable (e.g. ``"k_bit_off"``), each element's bit position is computed as ``bit_off_var + i * bits`` at runtime, still fully unrolled (one expression per element, no loops). z[kb[z] & zu]z[(kb[z] >> z) & z[((kb[z u) | ((kb[zu) << z)](r)z >> 3)z & 7)z[((z <= 8) ? (kb[z ) : ((kb[z) | (kb[z + 1] << (8 - z)))) & )rOr)rrrrrexprsi bit_offsetbyte_idx bit_in_byteexprlow_bits high_maskbobybis&&&& r_gen_unrolled_extractrs I?DE wAJ!QH$q.K!Q&!#+_D $tfBGD)/xjk]$tfTVW{?4?3q8 $oVH:U;-tQRZ]^_L_K`a'!|nD 6(2O LL ' D LwA451[MQXJa0;BRDBRDB /"DD6"U2$'$eB4xt>"FV2   LL   Lrc(V^8dQhRRRRRRRR/#r rr0rrrr r)rs"rrrs(..c.C.c.3.rc\WRV4p\V4UUu.uF wrERV RV 2NK pppRPV4#uuppi)z9Generate score accumulation with unrolled key extraction.rXzq[z] * z + )rrrP)rrrrrr termss&&& r_gen_unrolled_scorersQ !$ ME/8/? @/?GAr!D /?E @ ! & &u -- AsAc(V^8dQhRRRRRRRR/#rr)rs"rrrs(cCc3rc \WRV4pVUu.uFqDPRR4NK pp.p\V4F wrgVPRV RV RV R24K" RP V4#uupi) z5Generate value accumulation with unrolled extraction.rzkb[zvb[z o[z] = o[z] * factor + exp_score * z * vn;r?)rreplacerrrP)rrrrerQrr s&&& r_gen_unrolled_valuersz !$ ME.3 4eYYue $eE 4 EU# ~aSqc1J4&PVWX$ 99U  5sA0c$V^8dQhRRRRRR/#r key_bitsr0val_bitsr_r)rs"rrrs&CCsCcCCrc6\4'dV^8:gV^8:dR#V^ 8gV^ ,^8wdR#^V,^, p^V,^, pV^ ,pRT RT RT RT RWP,^,'dRMR R WQ,^,'dR MR R \YWP,^,'dR MR4 R \YWQ,^,'dRMR4 R2p\PP RV RV RV 2.ROR.VR7#)uKFused MSE decode: 32 simdgroups × 32 lanes, online softmax + weighted sum.Nz constexpr int BN = 32; // simdgroups per threadgroup constexpr int BD = 32; // lanes per simdgroup constexpr int qk_per_thread = Dim / BD; constexpr int v_per_thread = Dim / BD; constexpr uint k_mask = #u; constexpr uint v_mask = "u; constexpr int k_bits = !; constexpr int v_bits = u; typedef float U; // Thread identity — grid is total threads, so use threadgroup position auto bqh = threadgroup_position_in_grid.x; // batch*q_head index auto simd_gid = simdgroup_index_in_threadgroup; auto simd_lid = thread_index_in_simdgroup; auto token_count = key_norms_shape[2]; auto kv_heads = key_norms_shape[1]; auto bh = bqh / RepeatCount; // map q_head -> kv_head auto k_nm = key_norms + bh * token_count; auto k_pk = key_packed + bh * token_count * KPackedWidth; auto v_nm = val_norms + bh * token_count; auto v_pk = val_packed + bh * token_count * VPackedWidth; // Shared memory for cross-simdgroup reduction threadgroup U max_scores[BN]; threadgroup U sum_exp_scores[BN]; threadgroup U shared[BN * BD]; // Preload pre-rotated query into registers thread U q[qk_per_thread]; auto qr = queries + bqh * Dim + simd_lid * qk_per_thread; for (int i = 0; i < qk_per_thread; i++) q[i] = static_cast(qr[i]); // Initialize accumulators thread U o[v_per_thread] = {}; U max_score = -INFINITY; U sum_exp_score = 0; // Byte/bit offset for this lane's first element (constant across all tokens) int k_bit_start = simd_lid * qk_per_thread * k_bits; int v_bit_start = simd_lid * v_per_thread * v_bits; int k_byte_base = k_bit_start >> 3; int v_byte_base = v_bit_start >> 3; int k_bit_off = k_bit_start & 7;r2 int v_bit_off = v_bit_start & 7;u^ // KV loop: each simdgroup handles tokens simd_gid, simd_gid+32, ... for (int t = simd_gid; t < (int)token_count; t += BN) { U kn = static_cast(k_nm[t]); // Key score — unrolled byte extraction auto kb = (const device uint8_t*)(k_pk + t * KPackedWidth) + k_byte_base; U score = k_bit_offu; score = simd_sum(score) * kn; // Online softmax auto vb = (const device uint8_t*)(v_pk + t * VPackedWidth) + v_byte_base; U vn = static_cast(v_nm[t]); U new_max = max(max_score, score); U factor = fast::exp(max_score - new_max); U exp_score = fast::exp(score - new_max); max_score = new_max; sum_exp_score = sum_exp_score * factor + exp_score; // Value accumulation — unrolled byte extraction v_bit_offu} } // Cross-simdgroup reduction (matches MLX SDPA pattern) // 1. Communicate max and sum_exp across simdgroups if (simd_lid == 0) { max_scores[simd_gid] = max_score; sum_exp_scores[simd_gid] = sum_exp_score; } threadgroup_barrier(mem_flags::mem_threadgroup); // 2. Reduce max and sum across all simdgroups U sg_max = max_scores[simd_lid]; U new_max = simd_max(sg_max); U factor = fast::exp(sg_max - new_max); U total_sum = simd_sum(sum_exp_scores[simd_lid] * factor); // 3. Rescale this simdgroup's factor for the global max U my_factor = fast::exp(max_score - new_max); // 4. Transpose-reduce outputs through shared memory for (int i = 0; i < v_per_thread; i++) { shared[simd_lid * BD + simd_gid] = o[i] * my_factor; threadgroup_barrier(mem_flags::mem_threadgroup); o[i] = simd_sum(shared[simd_gid * BD + simd_lid]); o[i] = total_sum > 0 ? o[i] / total_sum : 0; threadgroup_barrier(mem_flags::mem_threadgroup); } // Write output (in rotated space — caller applies inverse rotation) if (simd_lid == 0) { for (int i = 0; i < v_per_thread; i++) { out[bqh * Dim + simd_gid * v_per_thread + i] = static_cast(o[i]); } } turboquant_fused_mse_sdpa_k_v_drrqueriesrQ key_packedrXrrrrrrrr)r*rrr_k_maskv_maskelems_per_laner"s&&& r_fused_mse_decode_kernelr4s[   Q(a- Rx38q=8mq F8mq FBYN! "()!!') (z* (z'* N1?0IQ/N/N +TVWX 0>0IQ/N/N +TVWX+8UcUnrsTsTs[y{|}~ !>KdhiJiJi;oq r s#tMiFV 77  *8*Bxj3% H W  rc$V^8dQhRRRRRR/#rr)rs"rrr s&ddsdcddrc8\4'dV^8:gV^8:dR#V^ 8gV^ ,^8wdR#^V,^, p^V,^, pV^ ,pRT RT RT RT RWP,^,'dRMR R WQ,^,'dR MR R \YWP,^,'dR MR4 R \YWQ,^,'dRMR4 R2p\PP RV RV RV 2.RO.ROVR7#)z92-pass decode pass 1: block-parallel quantized attention.Nz constexpr int BD = 32; constexpr int qk_per_thread = Dim / BD; constexpr int v_per_thread = Dim / BD; constexpr uint k_mask = r!r"r#u; typedef float U; // Thread identity — matches sdpa_vector_2pass_1 layout auto kv_head_idx = threadgroup_position_in_grid.x; auto batch_idx = threadgroup_position_in_grid.y; auto block_idx = threadgroup_position_in_grid.z; auto simd_lid = thread_index_in_simdgroup; auto gqa_idx = thread_position_in_threadgroup.y; // which q_head within kv_head auto token_count = key_norms_shape[2]; auto kv_heads = key_norms_shape[1]; auto bh = batch_idx * kv_heads + kv_head_idx; auto bqh = batch_idx * kv_heads * RepeatCount + kv_head_idx * RepeatCount + gqa_idx; auto k_nm = key_norms + bh * token_count; auto k_pk = key_packed + bh * token_count * KPackedWidth; auto v_nm = val_norms + bh * token_count; auto v_pk = val_packed + bh * token_count * VPackedWidth; // Load pre-rotated query thread U q[qk_per_thread]; auto qr = queries + bqh * Dim + simd_lid * qk_per_thread; for (int i = 0; i < qk_per_thread; i++) q[i] = static_cast(qr[i]); thread U o[v_per_thread] = {}; U max_score = -INFINITY; U sum_exp_score = 0; // Byte/bit offset for this lane's first element (constant across all tokens) int k_bit_start = simd_lid * qk_per_thread * k_bits; int v_bit_start = simd_lid * v_per_thread * v_bits; int k_byte_base = k_bit_start >> 3; int v_byte_base = v_bit_start >> 3; r$r2r%r&uc // KV loop: stride by blocks (each block handles a different subset) for (int t = block_idx; t < (int)token_count; t += Blocks) { U kn = static_cast(k_nm[t]); // Key score — unrolled byte extraction auto kb = (const device uint8_t*)(k_pk + t * KPackedWidth) + k_byte_base; U score = r'a; score = simd_sum(score) * kn; auto vb = (const device uint8_t*)(v_pk + t * VPackedWidth) + v_byte_base; U vn = static_cast(v_nm[t]); U new_max = max(max_score, score); U factor = fast::exp(max_score - new_max); U exp_score = fast::exp(score - new_max); max_score = new_max; sum_exp_score = sum_exp_score * factor + exp_score; r(a } // Write partial results for this block if (simd_lid == 0) { out_sums[bqh * Blocks + block_idx] = sum_exp_score; out_maxs[bqh * Blocks + block_idx] = max_score; } for (int i = 0; i < v_per_thread; i++) out_acc[(bqh * Blocks + block_idx) * Dim + simd_lid * v_per_thread + i] = static_cast(o[i]); turboquant_mse_sdpa_2pass1_kr*r+rr,)rout_sumsout_maxsr/r0s&&& r _fused_mse_decode_2pass_1_kernelr: sY   Q(a- Rx38q=8mq F8mq FBYN!"()!!') (z* (z#* F1?0IQ/N/N +TVWX 0>0IQ/N/N +TVWX+8UcUnrsTsTs[y{|} ~ !>KdhiJiJi;oq r s tJFX 77  +H:RzC5 I 9  rcv\4'gR#Rp\PPR.ROR.VR7#)z32-pass decode pass 2: reduce partial block results.Na constexpr int BN = 32; constexpr int BD = 32; constexpr int elem_per_thread = Dim / BD; typedef float U; thread U o[elem_per_thread] = {}; threadgroup U outputs[BN * BD]; auto head_idx = threadgroup_position_in_grid.x; auto simd_gid = simdgroup_index_in_threadgroup; auto simd_lid = thread_index_in_simdgroup; auto p = partials + head_idx * Blocks * Dim + simd_gid * Dim + simd_lid * elem_per_thread; auto s = sums + head_idx * Blocks; auto m = maxs + head_idx * Blocks; U max_score = -INFINITY; U sum_exp_score = 0; // Each simdgroup handles a subset of blocks (stride BN) for (int b = simd_gid; b < Blocks; b += BN) { U block_max = m[b]; U block_sum = s[b]; U new_max = max(max_score, block_max); U factor = fast::exp(max_score - new_max); U block_factor = fast::exp(block_max - new_max); sum_exp_score = sum_exp_score * factor + block_sum * block_factor; for (int i = 0; i < elem_per_thread; i++) { o[i] = o[i] * factor + partials[(head_idx * Blocks + b) * Dim + simd_lid * elem_per_thread + i] * block_factor; } max_score = new_max; } // Cross-simdgroup reduction (same as sdpa_vector_2pass_2) threadgroup U sg_maxs[BN]; threadgroup U sg_sums[BN]; if (simd_lid == 0) { sg_maxs[simd_gid] = max_score; sg_sums[simd_gid] = sum_exp_score; } threadgroup_barrier(mem_flags::mem_threadgroup); U sg_max = sg_maxs[simd_lid]; U new_max = simd_max(sg_max); U factor = fast::exp(sg_max - new_max); U total_sum = simd_sum(sg_sums[simd_lid] * factor); U my_factor = fast::exp(max_score - new_max); for (int i = 0; i < elem_per_thread; i++) { outputs[simd_lid * BD + simd_gid] = o[i] * my_factor; threadgroup_barrier(mem_flags::mem_threadgroup); o[i] = simd_sum(outputs[simd_gid * BD + simd_lid]); o[i] = total_sum > 0 ? o[i] / total_sum : 0; threadgroup_barrier(mem_flags::mem_threadgroup); } if (simd_lid == 0) { for (int i = 0; i < elem_per_thread; i++) out[head_idx * Dim + simd_gid * elem_per_thread + i] = static_cast(o[i]); } turboquant_mse_sdpa_2pass2rr)partialssumsmaxsr(r+s r _fused_mse_decode_2pass_2_kernelr@ sE   DFL 77   )0W  rc(V^8dQhRRRRRRRR/#r shared_namer sign_name temp_namedims_per_lane_paddedr)rs"rrr@ s0!!!!! !rc~RRV R2RV RV R2RRR R V R 2R R RV R2RV R2RV R2RRRR V R2RV RV R2RRRRR.#)zGenerate Metal code for in-place RHT forward in shared memory. RHT forward = WHT(signs * x) / sqrt(DimPadded). Requires: shared[DimPadded], sign_vec[Dim], temp[DimsPerLanePadded] (thread-local). Returns list of Metal source lines.z9 // RHT forward: apply signs, butterfly WHT, scale& for (int i = 0, d = lane; i < ; d += 32, i++) if (d < Dim) z[d] *= [d];rr2? for (int stride = 1; stride < DimPadded; stride *= 2) {* for (int i = 0, d = lane; i < ; d += 32, i++) {$ if (d < DimPadded) {* int pair = d ^ stride; float a = [min(d, pair)]; float b = [max(d, pair)]; %[i] = (d < pair) ? (a + b) : (a - b); }r8< threadgroup_barrier(mem_flags::mem_threadgroup);# if (d < DimPadded) [d] = rr4z% // Scale by 1/sqrt(DimPadded)8 float rht_scale = 1.0f / sqrt((float)DimPadded);rrCrDrErFs&&&&r_metal_butterfly_wht_forwardr]@ s D 01E0FoV #K= {$G B J 45I4JJ\] / 4 ( _E ( _E yk)NO   F 45I4J/Z -k]& 4P F  / B- rc(V^8dQhRRRRRRRR/#rBr)rs"rrrd s0 rc|RRRV R2RRRV R2R V R 2R V R 2R RRRV R2RV RV R2RRRRRRV R2RV RV R2R.#)zGenerate Metal code for in-place RHT inverse in shared memory. RHT inverse = signs * WHT(x) / sqrt(DimPadded). Same butterfly as forward, then multiply by signs.zA // RHT inverse: butterfly WHT, then apply signs and scalerLrMrNrOrPrQrRrSrTrUrVrWr8rXrIrYrZrr4r2z // Scale and apply signsr[rHrJz[d] *= rht_scale * rKrrr\s&&&&rrrd s L J 45I4JJ\] / 4 ( _E ( _E yk)NO   F 45I4J/Z -k]& 4P F  * B 01E0FoV #K=0CI;dS B+ rc V^8dQhRRRR/#)r rr0rr)rs"rrr syyysyrc F\4'dV^8:gV^8:dR#^V,^, p^V,^, p^V,^, p^V,^, pRV RV RV RV RV RV R2 p\PPRV R V 2.R O.R OVR 7#) z'Fused key+value quantize in 1 dispatch.Nah auto d = thread_position_in_threadgroup.x; auto bh = threadgroup_position_in_grid.x; auto is_val = threadgroup_position_in_grid.y; // 0=key, 1=value auto sg_id = simdgroup_index_in_threadgroup; auto sg_lid = thread_index_in_simdgroup; // Select key or value params based on grid.y int bits = is_val ? z : z; int n_mid = is_val ? z#; uint idx_mask = is_val ? zu : u u; int pw = is_val ? VPackedWidth : KPackedWidth; // Step 1: Load vector element float v; if (is_val) v = (d < Dim) ? static_cast(val_vectors[bh * Dim + d]) : 0.0f; else v = (d < Dim) ? static_cast(key_vectors[bh * Dim + d]) : 0.0f; // Compute norm (2-stage cross-simdgroup reduction) float sq = v * v; float sg_sum = simd_sum(sq); threadgroup float sg_norms[8]; if (sg_lid == 0) sg_norms[sg_id] = sg_sum; threadgroup_barrier(mem_flags::mem_threadgroup); float total_sq = (sg_id == 0 && sg_lid < 8) ? sg_norms[sg_lid] : 0.0f; total_sq = simd_sum(total_sq); if (sg_id == 0 && sg_lid == 0) sg_norms[0] = total_sq; threadgroup_barrier(mem_flags::mem_threadgroup); float norm = sqrt(sg_norms[0]); float inv_norm = (norm > 1e-10f) ? (1.0f / norm) : 0.0f; if (d == 0) { if (is_val) out_val_norms[bh] = half(norm); else out_key_norms[bh] = half(norm); } // Step 2: Unit vector → shared threadgroup float shared[Dim]; if (d < Dim) shared[d] = v * inv_norm; threadgroup_barrier(mem_flags::mem_threadgroup); // Step 3: Rotate (1 dim per thread, 256 FMAs) float rotated = 0.0f; if (d < Dim) { auto row = is_val ? (val_rotation + d * Dim) : (key_rotation + d * Dim); for (int j = 0; j < (int)Dim; j++) rotated += shared[j] * row[j]; } // Step 4: Comparison-based quantize threadgroup uint shared_idx[Dim]; uint idx = 0; if (d < Dim) { if (is_val) { for (int m = 0; m < n_mid; m++) idx += (rotated > val_midpoints[m]) ? 1u : 0u; } else { for (int m = 0; m < n_mid; m++) idx += (rotated > key_midpoints[m]) ? 1u : 0u; } shared_idx[d] = idx; } threadgroup_barrier(mem_flags::mem_threadgroup); // Step 5: Pack indices — thread-per-word, race-free, no atomics. // Each thread d (d < pw) walks the dims whose // [i*bits, (i+1)*bits) range intersects [32d, 32(d+1)) and // accumulates them into a private register, then writes the // word once. `bits` and `pw` are runtime-uniform here so the // loop bound is well-defined per dispatch. if (d < pw) { uint w_val = 0u; int word_start = (int)d * 32; int i_min = word_start / bits; int i_max = (word_start + 31) / bits; if (i_max >= (int)Dim) i_max = (int)Dim - 1; for (int i = i_min; i <= i_max; i++) { uint idx_val = shared_idx[i] & idx_mask; int bit_off = i * bits - word_start; if (bit_off >= 0) { w_val |= idx_val << bit_off; } else { w_val |= idx_val >> (-bit_off); } } if (is_val) out_val_packed[bh * pw + d] = w_val; else out_key_packed[bh * pw + d] = w_val; } turboquant_fused_kv_quantize_kr*r) key_vectors val_vectors key_rotation val_rotation key_midpoints val_midpoints) out_key_normsout_key_packed out_val_normsout_val_packedr()rr k_midpoints v_midpointsr1r2r"s&& r_fused_kv_quantize_kernelro s   Q(a-=A%K=A%K8mq F8mq F&Jc(4)]#k];""(fXR6\F| 77  -hZr( D  !  rcV^8dQhRR/#rMr)rs"rrr sIIsIrc \4'dV^8:dR#^V,^, p^V,^, pRV RV RV RV RV R2 p\PPRV 2R R .R .VR 7#) zQuantize pre-rotated vectors: comparison + pack. No rotation inside. Used with mx.hadamard_transform for external rotation. TG=Dim.Na auto d = thread_position_in_threadgroup.x; auto bh = threadgroup_position_in_grid.x; // Read pre-rotated, pre-normalized value float val = (d < Dim) ? static_cast(rotated[bh * Dim + d]) : 0.0f; // Comparison-based quantize threadgroup uint shared_idx[Dim]; uint idx = 0; if (d < Dim) { for (int m = 0; m < a; m++) idx += (val > midpoints[m]) ? 1u : 0u; shared_idx[d] = idx; } threadgroup_barrier(mem_flags::mem_threadgroup); // Pack indices: thread d builds word d in a private register // and writes it once (no shared-memory race, no atomics). // Walks every dim whose [i*bits, (i+1)*bits) range intersects // [32d, 32(d+1)), including a single spill from the previous // word for non-32-aligned bit-widths (3-bit, 5-bit, etc.). if (d < PackedWidth) { uint w_val = 0u; int word_start = (int)d * 32; int i_min = word_start / .; int i_max = (word_start + 31) / ; if (i_max >= (int)Dim) i_max = (int)Dim - 1; for (int i = i_min; i <= i_max; i++) { uint idx_val = shared_idx[i] & %u; int bit_off = i * a - word_start; if (bit_off >= 0) { w_val |= idx_val << bit_off; } else { w_val |= idx_val >> (-bit_off); } } out[bh * PackedWidth + d] = w_val; } turboquant_norot_quantize_rotated midpointsrrr()r num_midpointsrr"s& r_fused_norot_quantize_kernelry s   $Y!OM I?D( !"/0&'+V,--1F3015v6##'& )='FR 77  )$ 0 ,W  rc V^8dQhRRRR/#)r rr0rrr)rs"rrrP sXXSX4Xrc \4'dV^8:dR#^V,pV^, pV^, pRV RV RV RV RV R2 p\PPRV 2.R OR R .VR 7#) zBFused MSE quantize: norm + rotate + quantize + pack in 1 dispatch.Nu auto d = thread_position_in_threadgroup.x; // 0..Dim-1, one dim per thread auto bh = threadgroup_position_in_grid.x; auto sg_id = simdgroup_index_in_threadgroup; auto sg_lid = thread_index_in_simdgroup; auto vec_ptr = vectors + bh * Dim; // Step 1: Load & compute norm (2-stage cross-simdgroup reduction) float val = (d < Dim) ? static_cast(vec_ptr[d]) : 0.0f; float sq = val * val; float sg_sum = simd_sum(sq); threadgroup float sg_norms[8]; if (sg_lid == 0) sg_norms[sg_id] = sg_sum; threadgroup_barrier(mem_flags::mem_threadgroup); float total_sq = (sg_id == 0 && sg_lid < 8) ? sg_norms[sg_lid] : 0.0f; total_sq = simd_sum(total_sq); // Broadcast norm to all threads if (sg_id == 0 && sg_lid == 0) sg_norms[0] = total_sq; threadgroup_barrier(mem_flags::mem_threadgroup); float norm = sqrt(sg_norms[0]); float inv_norm = (norm > 1e-10f) ? (1.0f / norm) : 0.0f; if (d == 0) out_norms[bh] = half(norm); // Step 2: Unit vector → shared for rotation threadgroup float shared[Dim]; if (d < Dim) shared[d] = val * inv_norm; threadgroup_barrier(mem_flags::mem_threadgroup); // Step 3: Rotate — 1 dim per thread, 256 FMAs (was 2048 at TG=32) float rotated = 0.0f; if (d < Dim) { auto row = rotation + d * Dim; for (int j = 0; j < (int)Dim; j++) rotated += shared[j] * row[j]; } // Step 4: Comparison-based quantize threadgroup uint shared_idx[Dim]; uint idx = 0; if (d < Dim) { for (int m = 0; m < u; m++) idx += (rotated > midpoints[m]) ? 1u : 0u; shared_idx[d] = idx; } threadgroup_barrier(mem_flags::mem_threadgroup); // Step 5: Pack indices — thread-per-word, race-free, no atomics. // Each thread d (d < PackedWidth) walks the dims whose // [i*bits, (i+1)*bits) range intersects [32d, 32(d+1)) and // accumulates them into a private register, then writes the // word once. This avoids both the original `|=` race and the // GPU watchdog hangs that high-contention atomic_or sees on // some Metal devices. if (d < PackedWidth) { uint w_val = 0u; int word_start = (int)d * 32; int i_min = word_start / rrrsrta - word_start; if (bit_off >= 0) { w_val |= idx_val << bit_off; } else { w_val |= idx_val >> (-bit_off); } } out_packed[bh * PackedWidth + d] = w_val; } !turboquant_fused_mse_quantize_v2_ out_norms out_packedr)vectorsr8rwr()rr num_entriesrxrr"s&& r_fused_mse_quantize_kernelrP s   t)K!OM ?D+!V"/0& '+V,--1F3015v6##'& )AIFV 77  0 78!<0  rc V^8dQhRRRR/#)r rr0rrr)rs"rrr sTT#TTrc\4'dV^8:dR#^V,pV^, pV^, p.R@OpV'd-V.RAO, pV\RRRR 4, pV.RBO, pM V.RCO, pV.RNRNRNRNRNRNR NRV R2NRNR NRNRNRNRNRNRNRNRNRNRV R2NRV R2NRV R 2NR NR!NRNRNR"NR#NR$NR%NR&NR'NRNR(NR)NR*NR NR+NRNRNR,NRNR-NR NR.NR NR/NR NR0NRNRNRNR1NR2NRNR3NR4NR5NR NR6NRN, pV'dR7MRpV'dRMR8p\PP R9V V 2R:VR;RP V4R?7#)EzQFused Prod quantize: norm + rotate + MSE quantize + residual + QJL in 1 dispatch.Nr2rKr4rsharedrvrr3 if (d < Dim) {rr8z: // Step 4: MSE quantize + compute rotated residualz( float rot_residual[DimsPerLane];z- threadgroup uint shared_mse_idx[Dim];z uint idx = 0;z$ for (int m = 0; m < z; m++)zA idx += (rotated[i] > midpoints[m]) ? 1u : 0u;z+ float estimate = codebook[idx];z4 rot_residual[i] = rotated[i] - estimate;z1 if (d < Dim) shared_mse_idx[d] = idx;z# // Step 5: Pack MSE indicesz9 for (int w = lane; w < MsePackedWidth; w += 32) {z uint word = 0;z) for (int b = 0; b < 32; b += z) {z) int dim = (w * 32 + b) / r=z> if (dim < Dim) word |= (shared_mse_idx[dim] & zu) << b;z; out_mse_packed[bh * MsePackedWidth + w] = word;z // Step 6: Residual normz float res_sq = 0.0f;rz8 res_sq += rot_residual[i] * rot_residual[i];z0 float res_norm = sqrt(simd_sum(res_sq));z: if (lane == 0) out_res_norms[bh] = half(res_norm);zF // Step 7: QJL projection (rotated_residual @ combined_proj_t)z> // combined_proj_t = (mse_rotation @ projection_t).Tz; // so row d gives the projection for output dim dz5 if (d < Dim) shared[d] = rot_residual[i];z+ threadgroup uint shared_signs[Dim];z float proj = 0.0f;z5 auto row = combined_proj_t + d * Dim;z/ proj += shared[j] * row[j];zE shared_signs[d < Dim ? d : 0] = (proj >= 0.0f) ? 1u : 0u;z1 // Step 8: Pack sign bits (1 bit per dim)z: for (int w = lane; w < SignPackedWidth; w += 32) {z* for (int b = 0; b < 32; b++) {z% int dim = w * 32 + b;z> if (dim < Dim) word |= shared_signs[dim] << b;z7 out_signs[bh * SignPackedWidth + w] = word;r mse_rotationturboquant_fused_prod_quantize_rrwr'combined_proj_tr?r)z5 auto lane = thread_position_in_threadgroup.x;rz* auto vec_ptr = vectors + bh * Dim;r2z& // Step 1: Load & compute normz float v[DimsPerLane];z float partial_sq = 0.0f;rKzH float x = (d < Dim) ? static_cast(vec_ptr[d]) : 0.0f;z v[i] = x;z partial_sq += x * x;r4z0 float norm = sqrt(simd_sum(partial_sq));z@ float inv_norm = (norm > 1e-10f) ? (1.0f / norm) : 0.0f;z2 if (lane == 0) out_norms[bh] = half(norm);r2)z, threadgroup float shared[DimPadded];zF for (int i = 0, d = lane; i < DimsPerLanePadded; d += 32, i++)zI shared[d < Dim ? d : 0] = (d < Dim) ? v[i] * inv_norm : 0.0f;rz& shared[Dim + lane] = 0.0f;rr2)r2# float rotated[DimsPerLane];z@ for (int i = 0, d = lane; i < DimsPerLane; d += 32, i++)zB rotated[i] = (d < Dim) ? shared[d] * rht_scale : 0.0f;)z& threadgroup float shared[Dim];r3z5 if (d < Dim) shared[d] = v[i] * inv_norm;rr2rrKz float sum = 0.0f;rz2 auto row = mse_rotation + d * Dim;rz. sum += shared[j] * row[j];r8z rotated[i] = sum;r4)r}out_mse_packed out_res_norms out_signs)rr]rr)r*rP)rrrrxrrQrrs&& r_fused_prod_quantize_kernelr s   Qx-K!OMQH E&    - j#':        & > >D> 3> 8 > M > $ > &> /}oVD> L> > 6> ?> <> > C> !>" .#>$ D%>& %'>( 4H:TB)>* 4H:Q?+>, I RZ[->. />0 F1>2 3>4 5>6 +7>8 '9>: 8;>< C=>> ;?>@ EA>B C>D QE>F IG>H FI>J KK>L @M>N CO>P Q>R 6S>T MU>V )W>X %Y>Z @[>\ =]>^ :_>` a>b Pc>d e>f Cg>h i>j l Em>n %o>p 5q>r 0s>t Iu>v w>x By>z {>>E@$&K#*ZN 77  .xj F        yy  rc$V^8dQhRRRRRR/#)r r r0 high_bitsr/r)rs"rrrD s&QQQQCQrcH\4'dV^8dR#\V^, ^4p\V^, ^4p^V,^, p^V,^, pRp.RNRNRNRNRNRNRNR NR NR NR NR NRNRNRNRNRNRNRNRNRNRNRNRNRNRNRNRNRNRNRNRNRNRNR NR!NR"NR#NR$NR%NR&NR'V R(2NRNp\V4F,p VR)V R*V R+2R,V R-V R+2R.V R/V R02R1V R2V R02., pK. \V4Fp VR3V R42R5V R62R7V R62., pK VRR8., p\V4Fp VR9V R62., pK VR:R;RR?R@V RA2RBRCV RDV REV RF2RGV RH2RIRJ. , p\V4Fp VRKV RLV RMV RNV RO2 ., pK VRPRQ., pVR:RRRSRTRURVR@V RA2RBRCV RDV REV RF2RGV RH2RWRJ. , p\V4Fp VRKV RXV RYV RZV RO2 ., pK VRPRQ., p\V4Fp VR[V R\V R(2., pK V.R|O, p\V4FQp VR]V R^V RA2R_V R`V RaV R(2RbV RcV RdV R(2ReV RfV RdV R(2RgV RhV RiV RjV RA2 RkV RlV RiV RjV Rm2 ., pKS VRnR., pVRoV R(2., p\V4F'p VRpRqV RrV RA2RnRsRtV RA2RuV RA2RvV RA2Rn., pK) RwPV4p .R}Op \P P RxV RyV RzV 2V .R~OV R{7#)z?Fused SplitCodec decode: score + online-softmax + weighted-sum.NzDimLow + DimHighr@rhrAr2riz2 auto token_count = key_low_norms_shape[2];z/ auto kv_heads = key_low_norms_shape[1];rDrErFrlrmrnrGrHz; auto kl_norms = key_low_norms + base * token_count;zJ auto kl_mse = key_low_mse + base * token_count * KLMsePackedWidth;z= auto kl_res = key_low_res_norms + base * token_count;zO auto kl_signs = key_low_signs + base * token_count * KLSignPackedWidth;z< auto kh_norms = key_high_norms + base * token_count;zK auto kh_mse = key_high_mse + base * token_count * KHMsePackedWidth;z> auto kh_res = key_high_res_norms + base * token_count;zP auto kh_signs = key_high_signs + base * token_count * KHSignPackedWidth;z; auto vl_norms = val_low_norms + base * token_count;zM auto vl_packed = val_low_packed + base * token_count * VLPackedWidth;z< auto vh_norms = val_high_norms + base * token_count;zN auto vh_packed = val_high_packed + base * token_count * VHPackedWidth;z/ // Value bit offset for this lane's dimz+ bool is_low_val = val_dim < DimLow;zA int vd_local = is_low_val ? val_dim : (val_dim - DimLow);z2 int v_bits = is_low_val ? VLBits : VHBits;z% int v_bo = vd_local * v_bits;rurvz* uint v_mask = (1u << v_bits) - 1u;z. bool v_spills = (v_off + v_bits > 32);z" bool v_valid = val_dim < (r;z auto qrl_z% = q_rot_low + (base * RepeatCount + z ) * DimLow;z auto qpl_z& = q_proj_low + (base * RepeatCount + z auto qrh_z& = q_rot_high + (base * RepeatCount + z ) * DimHigh;z auto qph_z' = q_proj_high + (base * RepeatCount + r|r}r~r6rrLr5z {z; auto mse_t = kl_mse + t * KLMsePackedWidth;z? auto sign_t = kl_signs + t * KLSignPackedWidth;z; float kn = static_cast(kl_norms[t]);zR float ksr = kn * key_low_scale[0] * static_cast(kl_res[t]);z9 for (int d = lane; d < DimLow; d += 32) {z! int bo = d * r=z= uint idx = (mse_t[bo >> 5] >> (bo & 31));z% if (((bo & 31) + rrrz idx &= rz7 float code = key_low_codebook[idx];z@ uint sb = (sign_t[d >> 5] >> (d & 31)) & 1u;z ps_z += kn * static_cast(qrl_z1[d]) * code + ksr * (sb ? static_cast(qpl_z[d]) : -static_cast(qpl_rrWr8z; auto mse_t = kh_mse + t * KHMsePackedWidth;z? auto sign_t = kh_signs + t * KHSignPackedWidth;z; float kn = static_cast(kh_norms[t]);zS float ksr = kn * key_high_scale[0] * static_cast(kh_res[t]);z: for (int d = lane; d < DimHigh; d += 32) {z8 float code = key_high_codebook[idx];z += kn * static_cast(qrh_z1[d]) * code + ksr * (sb ? static_cast(qph_z[d]) : -static_cast(qph_r9r:rrrrrrrrrrrrrrrrrr4z int out_stride = (rrrrrrrr?turboquant_fused_split_decode_rrr)r2rrz! if (is_low_val) {z< auto vt = vl_packed + t * VLPackedWidth;4 uint vv = (vt[v_word] >> v_off);[ if (v_spills) vv |= vt[v_word + 1] << (v_bits - (v_off + v_bits - 32));z] v_code = val_low_codebook[vv & v_mask] * static_cast(vl_norms[t]);z } else {z< auto vt = vh_packed + t * VHPackedWidth;rrz^ v_code = val_high_codebook[vv & v_mask] * static_cast(vh_norms[t]);rWr8) q_rot_low q_proj_low q_rot_high q_proj_high key_low_norms key_low_msekey_low_res_norms key_low_signskey_high_norms key_high_msekey_high_res_normskey_high_signs val_low_normsval_low_packedval_high_normsval_high_packedkey_low_codebookkey_high_codebook key_low_scalekey_high_scaleval_low_codebookval_high_codebookr)rrFrOrPrr)r*) r rr/ low_mse_bits high_mse_bits low_mse_mask high_mse_mask val_dim_totalrQrRr"r s &&& r_fused_split_decode_kernelrC s   !1x!|Q'L A q)M%*L-'1,M&M, 8, <,  6,  ,  6 ,  ,  =,  :,  V,  /,  5,  *,  *,  2,  ,  8!, " L#, $ %, & F', ( U), * H+, , Z-, . G/, 0 V1, 2 I3, 4 [5, 6 7, 8 F9, : X;, < G=, > Y?, @ A, B :C, D 6E, F LG, H =I, J 0K, L *M, N )O, P 5Q, R 9S, T -]O2>U, V W, E\<  s"Gs+ Vs"H; Ws"H< Xs"I!L Y   !<  !!M 2!!H -!!H -  !  9E <  )!H566! EIE\ D +L>; G / ~=efrestDEQDRRZ [ %l^26 A J  E<  %aS(HK|}~|@_`a_bbh i  ! E  EIE] E +M?!< G / >fgtfuvFGTFUU] ^ %m_B7 B J  E<  %aS(HK|}~|@_`a_bbh i  ! E <  (?1#R@AA! E$<  #A3hqc 3s,qcaS ;#A3j8A3b A"1#Yqc!B ?s(1#VA3eA3a @s(1#VA3eA3j I   ! k2 E *=/< ==E <  %STUSVVwxywzz{ |*XYZX[[\ ]21#Q 721#Q 7  !YYu FK2 77  -hZq 2l^ T6  rc V^8dQhRRRR/#)r r r0rr)rs"rrr sppCpCprcaaaa\S^, ^4o\S^, ^4o\PRVVVV3Rll4pV#)z0Compiled SplitCodec decode handling both halves.c6V^8dQhRRRRRRRRRRRRRRR RR RR RR RR RRRRRRRRRRRRRRRRRRRRRRRRRRRRRRR/#)r rPrrkey_low_mse_indiceskey_low_residual_normskey_low_qjl_signsrkey_high_mse_indiceskey_high_residual_normskey_high_qjl_signsvalue_low_normsvalue_low_indicesvalue_high_normsvalue_high_indiceskey_low_transform_tkey_high_transform_trrrrvalue_low_codebookvalue_high_codebookvalue_low_rotationvalue_high_rotationlow_idxhigh_idx restore_orderr r)rs"rr3_compiled_split_decode_kernel..__annotate__ s9h7h7!h7 h7& h7 !) h7 $ h7!h7'h7"*h7%h7"h7$h7 #!h7"%#h7&&'h7(')h7*#+h7,$-h7. /h70!1h72%3h74&5h76%7h78&9h7<=h7>?h7@ Ah7B Ch7rc <VPR,p\P!VVRR7p\P!VVRR7p\P!VV 4pVPR,p\ VRRV13,P VP^,VP^,VP^,V4VRVR13,P VP^,VP^,VP^,V4\ VVVV4S*VV4p\P!VV4p VPR,p!\ V RRV!13,P V P^,V P^,V P^,V!4V RV!R13,P V P^,V P^,V P^,V!4\ VVVV4S(VV4p"VV",p#\V#\W4S)VV4p$\V#\W4S'VV4p%\P!V$V%.RR7p&\P!V&VRR7#)rnr.Nr) rrtaker;rr@rrHr concatenate)+rPrrrrrrrrrrrrrrrrrrrrrrrrrr_q_lowq_highqt_lowd_low scores_lowqt_highd_high scores_highrout_lowout_highmergedrrr rs+&&&&&&&&&&&&&&&&&&&&&&&&&& rr`._compiled_split_decode_kernel.._decode s>F##B'r:(<5"56 B& 3;  ' ' Qa&,,q/5  3;  ' ' Qa&,,q/5  #&!      &))F$89b!' C&L ! ) ) a '--"2GMM!4Df  CL ! ) ) a '--"2GMM!4Df  $'"      &k)6   B     7  / D     ( 3"=wwv}266rrb)r rr`rrsff @@r_compiled_split_decode_kernelr sKx!|Q'L A q)MZZh7h7h7T Nrc,]tRtRt$R]R&R]R&RtR#)ri rr%rrN__name__ __module__ __qualname____firstlineno____annotations____static_attributes__rrrrr s O rrc@]tRtRt$R]R&R]R&R]R&R]R&RtR#) ri rr%rr^rrNrrrrrr s Orrc,]tRtRt$R]R&R]R&RtR#)TurboQuantPolarStatei rrrr*rNrrrrrr s O''rrc@]tRtRt$R]R&R]R&R]R&R]R&RtR #) ri rr%rr)r^rrNrrrrrr s O%%rrc,]tRtRt$R]R&R]R&RtR#)TurboQuantSplitStatei objectlowhighrNrrrrrr s K Lrrc V^8dQhRRRR/#)r rfloatr r)rs"rrr s   5 rc\V4pV^8d \R4h\V^,4^, p\P!WRR7'g\RV R24hV#)rnz!TurboQuant requires kv_bits >= 1.r abs_tolz=TurboQuant currently supports integer and .5 bit-widths, got .)r ValueErrorroundmathisclose)rroundeds& r_validate_bitsr sb ;D ax<==D1Ho!G <<t 4 4KD6QR S   Nrc$V^8dQhRRRRRR/#)r rzOptional[float]schemez Optional[str]r rr)rs"rrr s!==_=m=t=rcVfR#VR8XdR#\V4p\P!V\V4RR7'*#)NF turboquantTr r)rrrr)rrs&&rturboquant_enabledr s8 |  ;D||D%+t< <F 99<< DA A 88AHHRZZ( ))rc V^8dQhRRRR/#)r nr0r r)rs"rrr s55#5#5rcPV^8d^V^, P4,#^#r) bit_length)r%s&r_next_power_of_twor( s%()A1Q""$ $414rc V^8dQhRRRR/#rr)rs"rrr s####rc\V4#r)r(rs&r_rht_padded_dimr+ s c ""rc$V^8dQhRRRRRR/#r r)rs"rrr s!#SXrcDV^8:d'\P!R\PR7#\PP WR,,4pVP RR.VR7P\P4p\P!V4#)zCDeterministic random sign vector for Randomized Hadamard Transform.rr?)sizer) rrrrrrchoicerr)r_rr r_s&& r_rht_sign_vectorr2 sn axxxBJJ// ))  Tz 1 2C JJc{J - 4 4RZZ @E 88E?rc$V^8dQhRRRRRR/#r xrr_r r)rs"rrr s!   H  X  (  rclVP^,p\V4pW,pW28dBR.VP^, ,^W2, 3.,p\P!WE4p\P !VR\ P!V4, R7pW28dVRRV13,pV#)z+RHT forward: hadamard(signs * x) / sqrt(D).r.r`.Nrrr+rrpadhadamard_transformrsqrtr5r_r_D_paddedy pad_widths&& r _rht_forwardr@ s ++a.Cs#H A~H +8>/B.CC FF1  asTYYx-@'@AA~ c4C4iL Hrc$V^8dQhRRRRRR/#r4r)rs"rrrs!  H X ( rclVP^,p\V4pTpW28dBR.VP^, ,^W2, 3.,p\P!WE4p\P !VR\ P!V4, R7pW28dVRRV13,pWA,#)z+RHT inverse: signs * hadamard(x) / sqrt(D).r.r7.Nrr8r<s&& rr:r:s ++a.Cs#H A~H +8>/B.CC FF1  asTYYx-@'@AA~ c4C4iL 9rc$V^8dQhRRRRRR/#r r)rs"rrrs!//C/s/x/rclV^8:d'\P!R\PR7#\PP WR,,^,4pVP W3\PR7p\P!VP\P44#)rri r) rrrrrrrrr)r_rr r!s&& r_projection_matrixrEsu axxxbjj11 ))  Tz 1B 6 7C  #2:: >F 88FMM"**- ..rc$V^8dQhRRRRRR/#)r r np.ndarrayr_r0r r)rs"rrrs!JSZrc V^8:d\P!V4pM\P!V^, 4R\P!\P 4,, \P!V^, ^, 4, pW1^, ^, \P!\P !RV^,, RR44,,p\P!V\P!V4, 4pVP4pV^8Xd(\P!VR\V4, 4#W%, #)rn?r.gKH9N) r ones_likerlgammalogpiclipexprFsum full_liker)rr_pdf log_coefflog_pdfpdf_sums&& r _beta_pdfrVs axll4  KKa 3$''):#: :T[[#PQ'UV=W W a1} GGC$'M5$ /1   ffWrvvg./ggiG!|||D#D /22 =rc$V^8dQhRRRRRR/#)r r_r0rr rr)rs"rrr*s!$2$23$2c$2h$2rcV^8:d'\P!R\PR7#^V,pV^8:d?\P!R RV\PR7p\P !V4#\P!R R R\PR7p\ W@4p\P!V4p\P!V\PR7R,V, p\P!WvV4P\P4p\^d4EFgp\P!V^,\PR7p R V ^&RV R &RVRR VR,,,V ^R %VP4p \V4Fp W^, 8Xd$WIV ,8WIV ^,,8*,p M"WIV ,8WIV ^,,8,p W\,p V P^8XdKpV P4pV^8gK\P!WV ,,4V, W&K \P !\P"!W, 44R8dT pMT pEKj \P !VP\P44#) rrr.r rINrnNNrr0g !g !?r)rrrrlinspacerrVcumsumarangeinterprrOemptycopyr/rPrFabs)r_rr, centroidsrrjcdf quantilesr boundaries new_centroidsrrbucket_weights total_weights&& r _codebookri)s qyxxBJJ// $YF axKKc6D xx "" ;;{JRZZ HD"G ))G C64s:fDI )$/66rzzBI 3ZXXfqj ;  1  2)CR.9R="@A 1R!( vAQJ1 -$QU:K2KL1 -$AE9J2JK$]N""a')--/La#%66.:*E#F#U   66"&&23 4t ;%I ! ), 88I$$RZZ0 11rc$V^8dQhRRRRRR/#)r rrGrr0r r)rs"rrrQs!  : c j rcV^8:d\P!V4pM_^V^, ,^, p\P!\P!\P!RV,4RR4V4pVP 4pV^8Xd(\P !VR\V4, 4#W$, #)rn@Nr.)rrJpowerrNsinrPrQr)rrrRexponentrUs&& r_polar_angle_pdfrqQs zll4 %!)$)hhrwwrvvcDj13=xHggiG!|||D#D /22 =rc$V^8dQhRRRRRR/#)r rr0rr rr)rs"rrr^s!&2&2&2C&2H&2rcV^8:d'\P!R \PR7#^V,pV^8:dR\P,V, p\ P !V\ PR7V,VR, ,p\P!VP\ P44#\ P!R\P^, R, R\ PR7p\WP4p\ P!V4p\ P !V\ PR7R,V, p\ P!WV4P\ P4p\^d4EF|p \ P!V^,\ PR7p RV ^&\P^, V R &RVRR VR,,,V ^R %VP4p \V4Fp W^, 8Xd$WZV ,8WZV ^,,8*,p M"WZV ,8WZV ^,,8,p Wm,pVP ^8XdKpVP#4pV^8gK\ P"!WV ,,4V, W&K \ P$!\ P&!W, 44R8dT pMT pEK \P!VP\ P44#) rrrlr rYrIrmNrZrr)rrrrrMrr]rrr[rqr\r^rOr_r`r/rPrFra)rr level_countsteprbrrjrcrdrrerfrrrgrhs&& r_polar_angle_codebookrv]sP qyxxBJJ//t)K zdgg ,IIk$>$>$>$>rc <V^8Xd9\P!.VPRRO^N5\PR7#VP \P4pVPR,p\ W!4pVP RV34p\4pVeV!V.RV3RV3RV3.W4P^,^3\^ V4^^3VP^,V3.\P.R7^,pVP .VPRROVN54#\P!VP^,V3\PR7p\V4FpWq,pV^ ,p V^ ,p VRV 3;;,VRV3,V ,,uu&W,^ , p V ^8gK\VRV ^,3;;,VRV3,W, , ,uu&K VP .VPRROVN54#) rNrrLengthrrNNNr) rrruint32rrzr@rWrrO) rVrrx packed_widthflatr r&idxrword_idxoffsetspills && r _pack_lowbitrs qyxx/&,,s+/Q/ryyAA ]]299 %F \\" F .L >>2v, 'D "F 66" -  1 q1R.15 JJqM<8999+     ~~@ Sb 1@<@AA XXtzz!}l3299 EFV}Z #bq({tAsF|v55 " 19 1hl? #tAsF| 'E E # >>>2v||B/0 1D " $F 66" 2/ **Q-+" JJqM62399+      <&,,s"3 .s+(Q$G1::(Qs(*#Unsupported TurboQuant state type: ) isinstancerrrr%rrrr^rrrtupler+r*r _concat_stater)rrr TypeErrortype)lhsrhss&&rrrsc {  { #)**! NNCIIsyy1 : NNCKK5A >  #*++" NNCIIsyy1 : NNCOOS__=A F NNC..0B0BC! L NNCMM3==9 B   #+,,# NNCIIsyy1 : E(+C,=,=s?P?P(QE  (+C,=,=s?P?P(Q   #/00' NNCIIsyy1 : #//3?? ; NNC..0B0BC! L NNCMM3==9 B   #+,,# #''377 + #((CHH -   9$s)G HHrcV^8dQhRR/#)r endr0r)rs"rrrsKKSKrc RaVfR#\V\4'd8\VPRRS13,VPRRS1R3,4#\V\4'de\ VPRRS13,VP RRS1R3,VP RRS13,VPRRS1R3,4#\V\4'dk\VPRRS1R3,\;QJd%.V3RlVP4FNK 54#!V3RlVP444#\V\4'dc\VPRRS13,\VPS4VP RRS13,VPRRS1R3,4#\V\4'd5\\VP S4\VP"S44#\%R\'V4: 24h)N.r~c3><"TFqRRS1R3,xK R#5i).Nr~r)rrrs& rr_slice_state..s#G3F%TcT1 %%3Fr)rrr%rrrr^rrrrr*r _slice_stater)rrrrr)rrs&frrrs }%+,,!%++c4C4i"8%--TcTST :UVV%,--" KKTcT "   c4C4l +  dsd + OOC#qL )   %-..# KKTcT1 % EG53F3FGE  G53F3FG G  %122' KKTcT " **C 0  dsd + OOC#qL )   %-..# C ( S )   9$u+I JJrc V^8dQhRRRR/#)r rr0rr)rs"rrr#s K KS Ks Krc ^aaVfR#\V\4'd8\VPRSS13,VPRSS1R3,4#\V\4'de\ VPRSS13,VP RSS1R3,VP RSS13,VPRSS1R3,4#\V\4'dm\VPRSS1R3,\;QJd&.VV3RlVP4FNK 54#!VV3RlVP444#\V\4'dd\VPRSS13,\VPSS4VP RSS13,VPRSS1R3,4#\V\4'd7\\VP SS4\VP"SS44#\%R\'V4: 24h)N.r~c3><"TFqRSS1R3,xK R#5i).r~Nr)rrrrs& rr%_slice_state_range..5s$L8KuU3Y)**8Krr)rrr%rrrr^rrrrr*r_slice_state_ranger)rrrrr)rrrs&ffrrr#s }%+,,! KKU3Y ' MM#uSy!+ ,  %,--" KKU3Y '   c59a/ 0  eCi 0 OOCsA- .   %-..# KKU3Y) * EL8K8KLE  L8K8KL L  %122' KKU3Y ' u00% =  eCi 0 OOCsA- .   %-..# uyy% 5 uzz5# 6   9$u+I JJrcV^8dQhRR/#r r r0r)rs"rrrFs   C  rcDVf^#\V\4'd0\VP4\VP4,#\V\ 4'd\ RV44#\V\P4'd VP#^#)Nc38"TFp\V4xK R#5ir) _state_nbytes)rparts& rr _state_nbytes..Ls954=&&5r) rrrrrrrPrrnbytesrs&rrrFsq }%-..UYY'- *CCC%95999%""|| rcV^8dQhRR/#rr)rs"rrrRs K KC Krc,Vf^#\V\4'd\VP4#\V\4'dVP P ^,#\V\4'dVP P ^,#\V\4'dVPP ^,#\V\4'dVP P ^,#\R\V4: 24h)Nr) rr _state_lengthrrr%rrrrrrrrs&rrrRs }%-..UYY''%+,,{{  ##%,--{{  ##%-..{{  ##%122{{  ## 9$u+I JJrcV^8dQhRR/#)r rxr0r)rs"rrrbs9K9K9Krc  a\V\4'd\\P!.VPP R,OSN5VPP R7\P!.VPP R,OSNVPP R,N5VPP R74#\V\4'Edq\\P!.VPP R,OSN5VPP R7\P!.VPP R,OSNVPP R,N5VPP R7\P!.VPP R,OSN5VPP R7\P!.VPP R,OSNVPP R,N5VPP R74#\V\4'd\\P!.VPP R,OSNVPP R,N5VPP R7\;QJd%.V3RlVP4FNK 54#!V3RlVP444#\V\ 4'Ed\!\P!.VPP R,OSN5VPP R7\#VP$S4\P!.VPP R,OSN5VPP R7\P!.VPP R,OSNVPP R,N5VPP R74#\V\&4'd5\'\#VP(S4\#VP*S44#\-R\/V4: 24h)Nr Nrc3<"TFPp\P!.VPR,OSNVPR,N5VPR7xKR R#5i)rrNr)rrrr)rrrxs& rr'_allocate_state_like..sS 1E ?ekk"o?v?u{{2?++1sAArr)rrrrr%rrrrrr^rrrrr*r_allocate_state_liker)rrrrr)rrxs&frrrbs%+,,! HH5u{{((,5f5U[[=N=N O HHK%--%%b)K6K5==3F3Fr3JKmm))   %,--" HH5u{{((,5f5U[[=N=N O HHS%##))"-SvSu7H7H7N7Nr7RS''--  HH9%&&,,R09&9**00  HHO%//''+OVOU__5J5J25NOoo++   %-..# HHG%++##B'GG1B1B21FGkk''  E #00 E   #00    %122' HH5u{{((,5f5U[[=N=N O !2!2F ; HH9%&&,,R09&9**00  HHO%//''+OVOU__5J5J25NOoo++   %-..# F 3 V 4   9$u+I JJrcV^8dQhRR/#)r rr0r)rs"rrrsII#IrcVfR#V\V4,p\V\4'd<VPVPRW#13&VPVPRW#1R3&R#\V\ 4'duVPVPRW#13&VP VP RW#1R3&VPVPRW#13&VPVPRW#1R3&R#\V\4'dQVPVPRW#1R3&\VPVP4F wrEWTRW#1R3&K R#\V\4'dyVPVPRW#13&\VPVPV4VPVPRW#13&VPVPRW#1R3&R#\V\ 4'dE\VP"VP"V4\VP$VP$V4R#\'R\)V4: 24h)N.r~r)rrrr%rrrr^rrrr+r*r _write_stater)rrrrr)dstsrcrr dst_level src_levels&&& rrrs { -$ $C#)**$'II #uy.!), CA%&#*++$'II #uy.!-0__UY)*-0-?-?3 >*+.== c59a'(#+,,'*yy #uy!#$$'(9(93;L;L$M I+4c59a' (%N#/00$'II #uy.!S__coou=-0-?-?3 >*+.== c59a'(#+,,SWWcggu-SXXsxx/ 9$s)G HHrc aVfR#\V\4'd/\S!VP^4S!VP^44#\V\4'dS\ S!VP^4S!VP ^4S!VP ^4S!VP^44#\V\4'df\S!VP^4\;QJd%.V3RlVP4FNK 54#!V3RlVP444#\V\4'dV\S!VP^4\VPS4S!VP ^4S!VP^44#\V\4'd5\\VP S4\VP"S44#\%R\'V4: 24h)a-Apply *fn* to every ``mx.array`` leaf in a TurboQuant state NamedTuple. *fn* receives ``(array, ndim)`` where *ndim* is the number of dimensions the array has (3 for norm-like ``(B,H,T)``, 4 for index-like ``(B,H,T,P)``). The function must return an ``mx.array`` with the same number of dimensions. Nc36<"TFpS!V^4xK R#5irNr)rrfns& rr_map_state..s@,?5"UA,,,?sr)rrr%rrrr^rrrrr*r _map_stater)rrrrr)rrs&frrrs }%+,,!"U[[!"4b6JKK%,--" u{{A  u  ! $ u##Q ' u "   %-..# u{{A  E@E,?,?@E  @E,?,?@ @  %122' u{{A  u((" - u##Q ' u "   %-..# uyy" % uzz2 &   9$u+I JJrc naVfVfR#\V\4'dE\S!VPVP^4S!VPVP^44#\V\4'd\ S!VPVP^4S!VP VP ^4S!VP VP ^4S!VPVP^44#\V\4'd\S!VPVP^4\;QJd9.V3Rl\VPVP44FNK 54#!V3Rl\VPVP4444#\V\4'd\S!VPVP^4\VPVPS4S!VP VP ^4S!VPVP^44#\V\ 4'dK\!\VP"VP"S4\VP$VP$S44#\'R\)V4: 24h)zAApply *fn(a1, a2, ndim)* element-wise across two matching states.Nc3:<"TFwrS!W^4xK R#5irr)rl1l2rs& rr"_map_state_pair..sV.UFB"RQ--.Usr)rrr%rrrr^rrrrr+r*r_map_state_pairr)rrrrr)s1s2rs&&frrrs zbj"())! rxx1 %r"**bjj!'D  ")**" rxx1 % r~~r~~q 1 r  ""3"3Q 7 r||R\\1 -   "*++# rxx1 % EVc"2B2BBDTDT.UVE  Vc"2B2BBDTDT.UV V  ".//' rxx1 % BNNBNNB ? r  ""3"3Q 7 r||R\\1 -   "*++# BFFBFFB / BGGRWWb 1   9$r(F GGrcV^8dQhRR/#r batch_indicesrr)rs"rrrs???rc$a\VV3Rl4#)z.Select batch elements from a TurboQuant state.c<VS,#rr)arrs&&r_filter_state.. s Q}-=rr)rrs&fr _filter_staters e= >>rc V^8dQhRRRR/#)r leftr0rightr)rs"rrr s # #3 #s #rcJaaS^8Xd S^8XdV#VV3Rlp\W4#)z(Pad along the token dimension (index 2).c<V^8Xd\P!VRRSS3.4#\P!VRRSS3R.4#)r)rr9)rrrrs&&r_pad_pad_state_tokens.._padsD 1966!fftUm<= =66!fftUmVDE Err)rrrrs&ff r_pad_state_tokensr s( qyUaZ F e ""rc\WR4#)z;Concatenate two states along the batch dimension (index 0).c0\P!W.4#rr)a1a2rs&&&rr%_concat_state_batch..sx8Pr)r)rrs&&r_concat_state_batchrs 2#P QQrc$V^8dQhRRRRRR/#)r usedr0neededrur)rs"rrr s!   c  rcVfR#\V4pWB8dV#W#,^, V,V,p\W4pV^8d\V\W4^4V#r)rrrr)rrrrucapacity new_capacitygrowns&&&& r_reserve_state_capacityr s\ }U#H ]Q&4/47L  5E axUL5q9 Lrc]tRtRtRRltRRltRRltRR ltR R ltR R lt RRlt RRlt RRlt RRlt RRltRRltRRltRRltRtR#) _TurboQuantMSECodeci.c$V^8dQhRRRRRR/#r r_r0rrr)rs"rr _TurboQuantMSECodec.__annotate__/s!??C?s?#?rc pWnW nV^8;'dW^, ,^8HVnVP'd\W4VnMRVn\ W4VnV^8dVP P4M VP Vn\W4Vn V^8dXVPP^,^8d6VPRRVPR,,^, Vn R#\P!R\PR7Vn R#)rNrZrrr)r_rrr2r_r#r8 transposerrir'r _midpointsrrrselfr_rrs&&&&r__init___TurboQuantMSECodec.__init__/s Qw99C7O#9 <<<)#4DJDJ(3 7:Qw$--113DMM!#, !8 ++A.2#}}Sb1DMM"4EEJDO hht2::>DOrc V^8dQhRRRR/#r r5rr r)rs"rrrAs---h-rc VP'd\WP4#\P!WP 4#r)rr@r_rr;rrr5s&&r_rotate_forward#_TurboQuantMSECodec._rotate_forwardAs- <<<::. .yyOO,,rc V^8dQhRRRR/#rr)rs"rrrFs+++h+rc VP'd\WP4#\P!WP 4#r)rr:r_rr;r8rs&&r_rotate_inverse#_TurboQuantMSECodec._rotate_inverseFs- <<<::. .yyMM**rc V^8dQhRRRR/#)r unit_vectorsrr tuple[mx.array, mx.array]r)rs"rrrKs??$? "?rc VP^8Xdi\P!.VPRRO^N5\PR7\P!VP\P R73#VP V4p\P!VP\PR7p\VPP^,4F<pW2VPV,8P\P4,pK> \W0P4p\P!VPVP\P4^R7pWPPV43#rNrrr)rrrrrrrrOrrrrr'int32r)rrrvrmr&estimated_rotateds&& r_quantize_unit_with_estimate0_TurboQuantMSECodec._quantize_unit_with_estimateKs 99>6<--cr26A6biiH++2::>  &&|4((7== :t,,Q/0A4??1+=!= E Ebii PPG1gyy1GGDMM7>>"((3KRST++,=>>>rc V^8dQhRRRR/#)r rrr r)rs"rrr\s00800rc VP^8Xd9\P!.VPRRO^N5\PR7#VP V4p\P!VP\PR7p\ VPP^,4F<pW2VPV,8P\P4,pK> \W0P4#)uBQuantize without computing the estimate (no wasted D×D rotation).Nrr) rrrrrrrOrrr)rrrvrr s&& r_quantize_unit"_TurboQuantMSECodec._quantize_unit\s 99>889l00"59q9K K&&|4((7== :t,,Q/0A4??1+=!= E Ebii PPG1GYY//rc V^8dQhRRRR/#)r packed_indicesrr r)rs"rrrfs--x-H-rc VP^8XdC\P!.VPRROVPN5\P R7#\ WPVP4P\P4p\P!VPV^R7pVPV4#r) rrrrr_rrrrrr'r)rrrrvs&& r_dequantize_unit$_TurboQuantMSECodec._dequantize_unitfs 99>88Bn223B7BB"**U U DHHELLRXXV''$--q9##G,,rc V^8dQhRRRR/#)r rrr rr)rs"rrrnsD D D -?D rc  HVPR ,^8XEdVP^8EdVP'EdVPpVP R V4P \ P4pVP^,p\ PPVR R7pV\ P!VR ,\4, pVPV4pW P,^,^ ,p\VP4p V eV !WpP.RV3RV3.W$,^^3V^^3WH3.\ P.R7^,p VPRR p \!VP \ P"4P V 4V P!.V OVN5!4#VPR ,^8XEdVP^8EdVP'Eg\%VPRR7p V EelVPpVP'd \'V4MTp VP R V4pVP^,pV^,^ ,p W P,^,^ ,pVP'd VP(M VP*pRV3RV 3RV3.pVP'dVR V 3R V ^,^ ,3., pV !W>VP.VW$,^^3V^^3V3WH3.\ P"\ P.R7wrZVPRR p\!VP V4V P!.VOVN5!4#VP \ P4p\ PPVR R7pV\ P!VR ,\4, p\!VP \ P"4VP-V44#)r rNrrrFrr DimPaddedrr?r.N)rrrr_r@rrrrnormmaximum_EPSrryrrrfloat16rr+r_r8r)rrrrBHr%unitrvrr r&origr=r  rot_inputr orig_shape vectors_f32rs&& rquantize_TurboQuantMSECodec.quantizensa ==  !dii!m A??2q)00bhmltr?rr~r~NNr~)rrr@rr'rrr_rrrreinsumr%r)rr0r fast_scoresrrvdotss&&& rscore_prepared"_TurboQuantMSECodec.score_prepareds  ! !" % ** (($**1-$**1-$**1-$**2.     K&""  488DKKBHHU''$--q9yy,.>Hkk((45HIIIrc$V^8dQhRRRRRR/#)r r-rrrr r)rs"rrrs'IIXI.@IXIrc DVPVPV4V4#rr7r-rr-rs&&&rscore_TurboQuantMSECodec.score ""4#7#7#@%HHrc$V^8dQhRRRRRR/#)r rjrrrr r)rs"rrrs"22H25G2H2rc &VPR,^8Xd?\VVVPVPVPVP 4pVeV#\ VPVPVP4P\P4p\P!VPV^R7p\P!RVVPP\P4V4pVP!V4#)r rzbhmlt,bht,bhtd->bhmldr?)rrCrr'r8r_rrr_rrrrr4r%rr)rrjr fast_outputrrvr7s&&& r weighted_sum _TurboQuantMSECodec.weighted_sums ==  !1    K&""  488DKKBHHU''$--q9yy #  KK  rzz *   ##L11rc$V^8dQhRRRRRR/#)r rrrrr r)rs"rrrs) E E E'9 E  Erc \VVVPVPVPVP4pVeV#VP \ P!VRR7V4#)Nrr)rHrr'r8r_rCrsoftmax)rrrrBs&&& rweighted_sum_from_scores,_TurboQuantMSECodec.weighted_sum_from_scoress\:   II MM MM JJ    "   F!\P!\P!WR,, 4RR7pWVV3#\P!WR,, 4pVPWr4p\P!VRR7pWV3#)rnrrr?r) rrFrrr@rKrr'r8r_rPrOrC) rrrr max_scores_2drBdenomrjrBs &&& rweighted_sum_stats_from_scores2_TurboQuantMSECodec.weighted_sum_stats_from_scoressVVF, ;;!  R 0A 5&..  #  #  #M B    K&rvvf)/D&DEBO":55&&Y"778""72wR(j((r)rrr'r_r8rr_rN)rrrrrrrr rrr%r)r-r7r=rCrHrOrrrrrr.sS?$- + ?"0-D LP-J.I2. E))rrc^]tRtRtRRltRRltRRltRR ltR R ltR R lt Rt R#)_PolarQuantUnitCodecic$V^8dQhRRRRRR/#rr)rs"rr!_PolarQuantUnitCodec.__annotate__s&WWCWsW#Wrc \V4'g\RV R24hWnW n\ W4Vn\ VP 4Vn\W4Vn V^8dVPP4M VPVn \;QJd).R\VP ^R74FNK 5M"!R\VP ^R744Vn\;QJd.RVP4FNK 5M!RVP44Vn\;QJd%.RVP4FNK 5VnR#!RVP44VnR#)z2PolarQuant requires a power-of-two dimension, got rc3<"TFwr\W4xK R#5ir)rv)rrrs& rr0_PolarQuantUnitCodec.__init__..(s"% %H! "% 4 4%Hsrc3N"TFp\P!V4xK R#5ir)rcosrr's& rrrW,VAUXx 0 0AU#%c3N"TFp\P!V4xK R#5ir)rrorZs& rrrW-r[r\N)rrr_rr rrr,r#r8rrrrangle_codebooksrr rs&&&&rr_PolarQuantUnitCodec.__init__s$$DSEK  +C6$//* (3 7:Qw$--113DMM$u% %.ta%H% uu% %.ta%H%   %VAUAUV%%VAUAUVV%VAUAUV%%VAUAUVVrc$V^8dQhRRRRRR/#)r anglesrrr0r r)rs"rrrT/s!;;h;s;x;rc \VPV^, ,p\P!VR,V, 4pV^8Xd4\P!VR\P ,V, 4p\P !VRR7P\P4#)rnrlrrr) r^rraminimumrrMargminrr)rrarr'diffss&&& r_quantize_level$_PolarQuantUnitCodec._quantize_level/sr'' 2vi(834 A:JJusTWW}&=>EyyR(// ::rc V^8dQhRRRR/#r rrr rr)rs"rrrT6s0Drc VPP\P4p\ \ VP 4\ VP4\ VP4\ VP44Fwr4rVVPR,p\WCV4P\P4p\P!WX^R7p \P!Wh^R7p \P!W),W*,.RR7P.VPRROVPR,^,N54pK V#)rnrNr)rrrrr+reversedrr*rr rrrrstackr@) rrrrindices_packedr-r. angle_countrcosinessiness && r_dequantize_preconditioned/_PolarQuantUnitCodec._dequantize_preconditioned6s ""2::.:= T__ % U(( ) T__ % T__ % ; 6D)  ++b/K$^;GNNrxxXGggiq9GGGIQ7EHHeou}=BGOO8%++cr"8EKKOa$78E;  rc V^8dQhRRRR/#)r rrr z%tuple[TurboQuantPolarState, mx.array]r)rs"rrrTGs""$" ."rc \P!WP4pTp.p\VP^R7FwrgVP .VP RROVP R,^,N^N54p\P!VR,VR,4p V^8Xd8\P!V ^8V R\P,,V 4p VPW4p VP\W44\PPVRR7pK \!VP#V4\%V44p VP'V 4p \P!WP(4p W3#)rnrNrlrr).rn).r)rr;rrrr@rarctan2whererrMrfrrrrrrrrqr8)rr storage_dtypepreconditionedr packed_levelsrrpairsrarrapprox_preconditioned approx_units&&& rquantize_unit_with_estimate0_PolarQuantUnitCodec.quantize_unit_with_estimateGs7<A $T__A>KEMM"NEKK$4"Nekk"o6J"NA"NOEZZf uV}=Fz&1*fsTWW}.DfM**69G  g!< =IINN5rN2E?% LL ' -  !% ? ? Fii 5}}E !!rc V^8dQhRRRR/#rir)rs"rrrT^sPP%9PhPrc b\P!VPV4VP4#r)rr;rqr8)rrs&&rdequantize_unit$_PolarQuantUnitCodec.dequantize_unit^s"yy88?OOrc(V^8dQhRRRRRRRR/#)r r0rrrr%r r)rs"rrrTas2DD (D1EDNVD Drc  VPR,^8Xd\VPVP^,VP^,VP^,VPR,4\VV\P !V4\P !.VPO^N5\PR74VPVPVP4pVeV#VPV4p\P!RW4pWcP\P4R,,#)r rr2r?rr3)rr0r@rr zeros_likerrrrr rqr4rr)rr0rr%r5r{r6s&&&& rr7#_PolarQuantUnitCodec.score_preparedas  ! !" % *1 (($**1-$**1-$**1-$**2.  )MM%(HH.u{{.A.bii@  K"&"" $ ? ? Fyy,.>Vll2::./BCCCr) r^rrr_rr,r8rr N) rrrrrrfrqr}rr7rrrrrRrRs)W$;"".PDDrrRc^]tRtRtRRltRRltRRltRR ltR R ltR R lt Rt R#)_TurboQuantPolarProdCodeci~c$V^8dQhRRRRRR/#rr)rs"rr&_TurboQuantPolarProdCodec.__annotate__s& D DC Ds D# Drc WnW n\WV4Vn\ W^,4VnV^8dVP P 4M VP VnV^8d9\P!VPPVP.RR7M%\P!R\PR7Vn V^8d3\P!\P ^, 4V, MRVn\P$!VP".\PR7VnR#rnrrrmNrr)r_rrR polar_codecrE projectionr projection_trrrrrquery_transform_trr;rMr`r scale_arrayrs&&&&rr"_TurboQuantPolarProdCodec.__init__s /4@,S(;;>7DOO557Qw NND,,779J9JKRT U& 3  691WTYYtww{+c1# 88TZZL Crc V^8dQhRRRR/#)r rrr rr)rs"rrrs   -E rc  VP\P4p\PP VRR7pV\P !VR,\ 4, pVPPVVPR7wrVWF, p\PP VRR7p\P!WpP4p \P!V ^8^^4P\P4p \VP\P4VVP\P4\!V ^44#)rnr)rwrr)rrrrrrrrr}rr;rrvrrrr) rrr$r%rr)r|residualr^ projectedr_s && rr%"_TurboQuantPolarProdCodec.quantizesnnRZZ0  {4"RZZi0@$%GG #'#3#3#O#O !--$P$   -r:IIh(9(9: aA.55bii@' LL $   ! !"** -  "   rc V^8dQhRRRR/#)r rrr rr)rs"rrrs S S 8 SX Src .VPPVP4p\VP^VP 4P \P4pVR,R, pVPVPR,P \P4,\P!W@P4,pVPR,P \P4W%,,#rnrlr.r)rrr)rrr_rrrr`r^r;rr%)rr polar_unit sign_bitsr_qjl_units&& rr)$_TurboQuantPolarProdCodec.dequantizes%%55e6G6GH "5??Atxx@GG S C#% JJ""9-44RZZ@ Aii/ 0  {{9%,,RZZ8J/>/3>/(>/  >/rc  Vwr4VPR,^8XEd \VPVP^,VP^,VP^,VPR,4VPVP^,VP^,VP^,VPR,4VVPPVPP VPP VP4pVeV#VPPVVPVP4pVPR,^8Xdz\VPVP^,VP^,VP^,VPR,4VVP4pVe Wg,#\VP^VP4P\ P"4pVR,R, p VP$VP&P\ P"4R,,\ P(!RVV 4,p VPP\ P"4R,p WkV ,,#)r rlr.r2r?rr3)rr5r@rrrr rr7r)r%rrrr_rrrr`r^r4) rr0r polar_queries proj_queriesr5 polar_scorefast_qjlrr_ qjl_scorer%s &&& rr7(_TurboQuantPolarProdCodec.score_preparedsX '7#   b !Q &2%%!''*!''*!''*!''+  $$ &&q) &&q) &&q) &&r*    ++  ++  ++  #K&&""&&55     KK   b !Q &'$$ &&q) &&q) &&q) &&r*     H#"--"5??Atxx@GG S C#% JJ""))"**56IJ Kii#   ""2::./BCY...rc$V^8dQhRRRRRR/#)r r-rrrr r)rs"rrrs'IIXI.FI8Irc DVPVPV4V4#rr;r<s&&&rr=_TurboQuantPolarProdCodec.scorer?r)rr_rrrrr`rN rrrrrr%r)r-r7r=rrrrrr~s+ D * SJ>/@IIrrc^]tRtRtRRltRRltRRltRR ltR R ltR R lt Rt R#)_TurboQuantProdCodecic$V^8dQhRRRRRR/#rr)rs"rr!_TurboQuantProdCodec.__annotate__s&GGCGsG#Grc WnW n\V\V^, ^4V4Vn\ W^,4VnV^8dVP P4M VP VnV^8d9\P!VPPVP.RR7M%\P!R\PR7VnV^8d3\P !\P"^, 4V, MRVn\P&!VP$.\PR7VnV^8d=\P*!VP VPP4VnR#\P!R\PR7VnR#r)r_rrrF mse_codecrErrrrrrrrrrr;rMr`rrr;_combined_proj_trs&&&&rr_TurboQuantProdCodec.__init__s+ ,S#dQh2BDI,S(;;>7DOO557Qw NNDNN55t7H7HIPR S& 3  691WTYYtww{+c1# 88TZZL C 7$&II!:!:%D !%'HHV2::$FD !rc V^8dQhRRRR/#)r rrr rr)rs"rrr sK K K -@K rc  <VPPpVPPpVPR ,^8XEdV^8Ed^ V,^8XEd\ W#R7pVEeVP pV'd \ V4MTpVPR V4pVP^,pV^,^ ,p WR,^,^ ,p V^,^ ,p V'dVPPMVPPp RV3RV 3RV 3RV 3.p V'dV RV3RV^,^ ,3., p V!VV VPPVPPVP.V ^ ^V3R V3W3V3W3.\P\P\P\P.R 7wrppVPRR p\!VPV4VP!.VOV N5!VPV4VP!.VOV N5!4#VP#\P$4p\P&P)VR R 7pV\P*!VR,\,4, pVPP/V4wppVV, p\P&P)VR R 7p\P0!VVP24p\P4!V^8^^4P#\P4p\!VP#\P4VVP#\P4\7V^44#)r rNrrrrrrrrr?rrr)rrrrrr_r+r@r_r8rr'rrrrrrrrrrrr r;rrvr)rrrrr rr=rrr mse_packed_widthsign_packed_widthr"rr%re res_normsr_r#r$rrmse_unitrr^rs&& rr%_TurboQuantProdCodec.quantize s>>&&..(( ==  !hlrH}7I0KF!HH18?1-ar1-ZZ]!"RB $%L2$5"#< %&VN!4;DNN00AXAX AJ"M2%'78&(9:  $h/,x"}.CD!H7=!11//-- &a *./ #    ##739e0%]]3B/ *MM*-&&E E4DE%%j1MMA:A/@A nnRZZ0  {4"RZZi0@$%GG $ K K !  X (*r:IIh(9(9: aA.55bii@" LL $   ! !"** -  "   rc V^8dQhRRRR/#)r rrr rr)rs"rrrYs Q Q 3 Q Qrc .VPPVP4p\VP^VP 4P \P4pVR,R, pVPVPR,P \P4,\P!W@P4,pVPR,P \P4W%,,#r)rrrrrr_rrrr`r^r;rr%)rrrrr_rs&& rr)_TurboQuantProdCodec.dequantizeYs>>2253D3DE"5??Atxx@GG S C#% JJ""9-44RZZ@ Aii/ 0  {{9%,,RZZ8H> ! ! !NN227;EYYw(9(9:F= ii)?)?@ 3 $(( ?+[dhhj-IIIrc$V^8dQhRRRRRR/#)r r0rrrr rr)rs"rrrls-H-H-3H-#H-  H-rc  Vwr4VPR,^8Xd\VPVP^,VP^,VP^,VPR,4VPVP^,VP^,VP^,VPR,4VVPPVPP VP 4pVeV#VPP^8d<VPPV\VPVP44pM\P!VP^,VP^,VP^,VP^,VPP^,3\PR7pVPR,^8Xdz\VPVP^,VP^,VP^,VPR,4VVP 4pVe Wg,#\VP ^VP"4P%\P4pVR,R, p VP&VP(P%\P4R,,\P*!RVV 4,p VPP%\P4R,p WkV ,,#)r rrlr.r2r?rr3)rrr@rrr'rr7rr%rrrrrrrr_rr`r^r4) rr0r mse_queriesrr5 mse_scorerrr_rr%s &&& rr7#_TurboQuantProdCodec.score_preparedls %5!   b !Q &+##%%a(%%a(%%a(%%b)  $$ &&q) &&q) &&q) &&r*  ##''  !K$&"" >>   "55"5;;0A0ABI  &&q) &&q) &&q) &&q)KK%%a( jj I   b !Q &'$$ &&q) &&q) &&q) &&r*     H# ++"5??Atxx@GG S C#% JJ""))"**56IJ Kii#   ""2::./BC9,,,rc$V^8dQhRRRRRR/#)r r-rrrr r)rs"rrrs'IIXI.AIhIrc DVPVPV4V4#rr;r<s&&&rr=_TurboQuantProdCodec.scorer?r) rrr_rrrrr`rNrrrrrrs.G,K Z QJH-TIIrrc$V^8dQhRRRRRR/#)r tensorravg_bitsrr ztuple[np.ndarray, np.ndarray]r)rs"rrrs$  %"rcD\P!V4p\P!V4pW#8Xd \R4hVPR,p\ \ W, V,W2, , 44p\^\V^, V44p\P!\P!VP\P44RR7p\P!\P !V44p\P"!Wu)RP\P$44p\P&!V\(R7p RW&\P*!V 4^,P\P$4p W3#)z:Mixed-precision selection requires a fractional bit-width.rNrFr)rrnr )rfloorceilrrr0rrFrrmeanrarrrargsortasarraysortrrrnonzero) rr lower_bits upper_bitsr_ high_countrorderrlow_maskrs && r_select_outlier_indicesrsH%J8$JUVV ,,r CUH1S8JbhhO,0FN(@S "3w<$G#CM4??D2IN 6>j9MNN,,r"CWBXBxxa"fq2vo 6bjjIH*...*J*J1crc6*RH\\3B3& '26..2R2R233H\\2B;. /11!SbS&9 ]]AFQVb[$88 96:__5V5V236H]]AFRKM1 2/7 +.2D +rc V^8dQhRRRR/#)r rrr rr)rs"rrrs  x ,@ rc \P!WPRR7p\P!WPRR7p\ VP P V4VPP V44#rnrr)rrrrrrr%r)rr low_tensor high_tensors&& rr%_SplitCodec.quantizes[WWV\\; ggfmm"= # NN # #J / OO $ $[ 1  rc V^8dQhRRRR/#r rrr rr)rs"rrrs<< 4< NN * *: 6 OO + +K 8  rc V^8dQhRRRR/#rr)rs"rrr s((6J(x(rc Vwr4VPPW2P4pVPPWBP4pWV,#r)rr7rrr)rr0r low_queries high_queries low_scores high_scoress&&& rr7_SplitCodec.score_prepared sC$4! ^^22; J oo44\::N ''rc$V^8dQhRRRRRR/#)r r-rrrr r)rs"rrrs'IIXI.BIxIrc DVPVPV4V4#rr;r<s&&&rr=_SplitCodec.scorer?rc$V^8dQhRRRRRR/#)r rjrrrr r)rs"rrrs"<>vzzR  9Cwwv11;;rc$V^8dQhRRRRRR/#)r rrrrr rKr)rs"rrr%s) O O O'; O , Orc VPPWP4wr4pVPPWP4wpp\ P !W6.RR7p\ P!WPRR7WE3#r) rrOrrrrrrr) rrrrrNrrrrs &&& rrO*_SplitCodec.weighted_sum_stats_from_scores%sy)-(U(U II) % :!OOJJ JJ  Q 9Cwwv11;UNNr) rrr_rrrrrrrrN)rrrrrr%r)r-r7r=rCrHrOrrrrrrs8$3L <  (I< < O Orrc(V^8dQhRRRRRRRR/#rr)rs"rrr3s(1111c11rc\V4p\P!V\V4RR7'dAVR8Xd\M\ pV!VP R,\\V44V4#\WW#4#)r rrr) rrrrrrrr0r)rrrrrs&&&& r _build_codecr3s` $ D ||D%+t44,0FN(@S b)3uT{+;TBB vT 00rc6]tRtRtRtRtRRltRtRtRt R#) _QuantizedStateProxyi;zWraps a quantized state tuple, providing .shape for model compatibility. Some models access keys.shape[-2] after cache.update_and_fetch() to slice masks. This proxy makes that work without dequantization. c V^8dQhRRRR/#)r n_tokensr0n_headsr)rs"rr!_QuantizedStateProxy.__annotate__Ds///c/rc &Wn^W2^3VnR#)rnN_stater)rrrrs&&&&rr_QuantizedStateProxy.__init__Ds A. rc .\VPV4#r)getattrr)rrs&&r __getattr__ _QuantizedStateProxy.__getattr__Ist{{D))rc ,\VP4#r)iterrrs&r__iter___QuantizedStateProxy.__iter__LsDKK  rrN) rrrr__doc__ __slots__rrr#rrrrrr;s $I/ *!rrc]tRtRtR)tRt^tRt]3RRllt ] ]3RRll4t RR lt R t R R lt]R 4tR*RltRRltR+RRlltR+RRlltRRltRRltRRltR+RRlltRt]R4t]P8R 4t]R!4t]P8R"4tR#tR$tR%t R&t!]R'4t"R(t#R#),TurboQuantKVCacheiPc V^8dQhRRRR/#)r rrrr0r)rs"rrTurboQuantKVCache.__annotate__Vs # #U ## #rc \V4VnW n^VnRVnRVnRVnRVnRVnRVn RVn RVn R#)rNr) rrrrkeysrV key_codec value_codec _cached_state_cached_state_offset _shadow_keys_shadow_values)rrrs&&&rrTurboQuantKVCache.__init__VsW"4(     !$&! "rc$V^8dQhRRRRRR/#)r rrrr0r z'TurboQuantKVCache'r)rs"rrr,ds$'* rc ^V!W#R7pVPwrVVeVPWV4V#)rr)rupdate_and_fetch)clscacherr turbo_cacher.rVs&&&& r from_cacheTurboQuantKVCache.from_cachecs3t/ {{     ( ( 6rc V^8dQhRRRR/#r r.rrVr)rs"rrr,ms8Xrc dVPf\P!VP\ VP4RR7'g!\P !VP4M VPp\ WRVPR7VnVPf\P!VP\ VP4RR7'g!\P!VP4M VPp\ W$RVP^,R7VnR#R#)Nr rmse)rr) r/rrrrrrrr0r)rr.rVrrs&&& r_ensure_codecs TurboQuantKVCache._ensure_codecsms >> ! ||DIIuTYY/?NN 499%YY  *$u499UDN    #||DIIuTYY/?NN $))$YY  ,u499q= D  $rc  VPR,^8wgA\VP\4'd!\VP\4'gR#VPP 'gVPP 'dR#\ VPP4p\ VPP4p\W44pVfR#VPR,pVPRV4pVPRV4pVP^,p Wc,^,^ ,p Wd,^,^ ,p V!VVVPPVPPVPPVPP.RV3RV 3RV 3.Wi,^^3V^^3V 3W3V 3W3.\P\P\P\P.R7wrrVPRRp\V PV4V P!.VOV N5!4\VPV4VP!.VOV N5!43#) zWFused key+value quantize in 1 dispatch. Returns (key_state, val_state) or (None, None).Nr KPackedWidth VPackedWidthrr?NNr)rrr/rr0rr0rror@r8rrrrr)rr.rVrrr rk_flatv_flatrk_pwv_pwk_normsk_packedv_normsv_packedr!s&&& r_try_fused_kv_quantize(TurboQuantKVCache._try_fused_kv_quantizes( JJrNa dnn.ABBd..0CDD  >> ! ! !T%5%5%=%=%= t~~**+t'',,-*8> >  JJrNb!$A& \\!_ r!b( r!b(/5''  ))))  ++  && &!QAq    ::ryy"**biiH-0 ,72zz#2 wt4h6F6F6S6Sd6S T wt4h6F6F6S6Sd6S T  rc V^8dQhRRRR/#r@r)rs"rrr,s% % X% x% rc jVPW4VPW4wr4Vf7VPPV4pVPPV4pVP VP ^,,pVPf"\W54Vn\WE4Vn M`\VPVP WPP4Vn\VPVP WPP4Vn \VPW0P 4\VPW@P 4VP ^,VP ^,rvVP R,pVP ^,p WPnRVn RVnV ^8gVP ^2,^8Xd,\P !VPVP4VP"wr\%WP V4\%WP V43#Nr)rCrQr/r%r0rrr.rrVr cache_steprr1r2revalrr) rr.rVnew_keys new_valuesnew_endrrrn_newksvss &&& rr9"TurboQuantKVCache.update_and_fetchs D) $::4H  ~~..t4H))226:J++ 1 - 99 ,X?DI.zCDK/ 4;;DI2 T[['??DK TYY++6T[[*kk:ZZ]DJJqM7 JJrN 1  !$&! 19r)Q. GGDIIt{{ + [[' : [[' :  rc J\V\4'd VP#T#r)rrrrs&r_unwrapTurboQuantKVCache._unwraps)%1EFFu||QEQrNc XVeVfVPwrVPV4pVPV4pVPPV4P \ P 4pVPPV4P \ P 4pW43#r)rr`r/r)rrrr0)r keys_state values_stater.rVs&&& rr)TurboQuantKVCache.dequantizes  !5'+zz $J\\*- ||L1 ~~((4;;BJJG!!,,\:AA"**M|rc<V^8dQhRRRRRRRRRRR RR RR RR R/ #) r rrrrq_startr0q_endk_startk_end total_queries total_tokensr r)rs"rrr,sd###!# #  #  #### #rc  VfV#\V\4'dVR8XdW, p \P!W,W,4p \P!WV4p V R,V R,8p V R,p \P!W\P !VP 4P4#\RV 24hVRW41WV13,p V PVP^, 8Xd\P!V ^R7p V P \P8Xd@\P!W\P !VP 4P4#W,#)Ncausalz'Unsupported TurboQuant attention mask: .r)r~N)Nr~)NNNr~r~) rrrr]rvfinforrrrrbool_)rrrrgrhrirjrkrl past_tokensq_idxk_idx causal_mask mask_chunks&&&&&&&&& r_apply_attention_mask'TurboQuantKVCache._apply_attention_masks <M dC x*:  +"79LM '1#Gng> )*@A xx RXXfll5K5O5OPPFtfMN N#w}gm;< ??fkkAo - ;J   rxx '88J0F0J0JK K""rc(V^8dQhRRRRRRRR/#r r-rr`rrrr r)rs"rrr,s8],],],  ], ! ], ],rc  VeVfVPwr#VPV4pVPV4pVPwrgr\V\4'd(VP P P^,MVP P^,p Wz,p W,PVV V VV 4p VPPp \V4pV^8Xd VPM VPpV^8Xd^M VPp.p\^VV4EFnp\VVV,4pV RVV1R3,pVP P#V4p\$P&!WjV VV, V 3\$P(R7p\$P&!WjV VV, 3\$P(R7p\$P*!WjV VV, 3\-R4)\$P(R7p\^W4EF,p\VVV,4p\/VVV4p\/VVV4pVP P1VV4pVP3VVVVVVVV4pVPP5VV4wppp \$P6!VV 4p!\$P8!VV!, 4p"\$P8!V V!, 4p#VV"R,,VV#R,,,pVV",VV#,,pT!p\$P:!VVV4EK/ VP=V\$P6!VR,\>4, 4\$P:!VR,4EKq \$P@!V^R7pVPWgW4pVPCVPD4#)N.r~rinfrrr)#rr`rrrrr%r@r0r_rdecode_key_chunk_sizeprefill_key_chunk_sizeprefill_query_block_sizerOrr/r-rrrfullrrr7rvrOrrOrWrrrrr)$rr-rcrdr`rr n_q_headsLr n_kv_heads n_repeatsrP value_dimrlkey_chunk_sizequery_block_sizeoutputsrgrhq_blockr0rB normalizer max_scorerirj key_chunk value_chunkr chunk_output chunk_denom chunk_maxnew_max prev_scale chunk_scales$&&&&&& rquantized_attention%TurboQuantKVCache.quantized_attentionsf  !5'+zz $J\\*- ||L1 $]]a*&:;; NN & &q )!!''*  + "?33    $$(( $Z0 *+q&D & &d6Q6Q !"Q1D,I,IQ#34G7%556E%c75=!&;DHDHDH  DH ! DH DHrc  V VeVfVPwr#VPV4pVPV4p\VP\4'dM\VP \ 4'd-\V\4'd\V\4'gR#VPwrgrVPP^,p Wz,p VPP^,p V ^8XdR#\VP P4p WP P8wdR#V ^,^ ,p\WV,V4pVfR#\VPPPV VV4pVfR#W,P!WjWV 4p\"P$!VVPP&4pVRRV 13,P!Wj,V ,W4pVRV R13,P!Wj,V ,W4pV!VVVPVP(VP*VP,VPPP.VPP0.RV 3RVP(PR ,3RVP,PR ,3.^ WV ,V ,3RWj,V ,W3.\"P2.R7^,pVP!WjWV 4pVe\V\"P44'dVP6VP6^, 8Xd\"P8!V^R7pVP:\"P<8XdB\"P>!VV\"P@!VP:4PB4pM"VV,pM\V\D4'dR#\"PF!VR R7P!Wj,W,V 4pRpV V,^, V,pWj,V,W,V 3pV!VVPVPHVP P..RV 3R W,3R V3R V3R VPHPR ,3.V ^Wj,V,3V ^^3V.\"P2.R7^,pVP!Wj,VW,V 4pV^8d\"PJ!V^R7pMVPM^4pVP!WjWV 4pVP POV4pVP!WgW4PQVP:4#)uFast prefill: fold L queries into R dimension, reuse decode kernels. Avoids the expensive O(T×D²) dequantize rotation matmul.N.rKMsePackedWidthKSignPackedWidthrrrrrrrr))rr`rr/rr0rrrrr%r0rrrrr@rr;rrr^rr'rrrrrrrprvrorrrGrrPsqueezerr)rr-rcrdr`rrrrrrrrrr  val_kernelmq_scoregroupedqtr$r]rrjr r  out_shape out_tiled out_rotatedrBs&&&&&& rprefill_attention#TurboQuantKVCache.prefill_attentionms  !5'+zz $J\\*- ||L1  t~~'; < <4++-@AA:':;;<);<<$]]a%%++A. +    " "1 % 6t'',,- '',, ,RB ; !m]   2 NN $ $ ) )      ?++A9K YYw @ @ A37 ##ANY$>ECG$$Q^i%?F  &&))$$((11**  "J$:$:$@$@$DE#Z%9%9%?%?%CD iZ!!34"NY6=>::,' ( ) ,yQ?  $))99 a/>>$Q7D::)XXdFBHHV\\4J4N4NOF#d]FD#&&**V"-55animUVW ]*Q.=@ ^m3Y]AF ""$$  ))   . . . 4 4 : :2 >? Q67Aq $+::,# $ % *%%anmY]TUV 1 &&3K#++A.K!))!qI !!11+>~~aA188GGrc V^8dQhRRRR/#r rPrr rr)rs"rrr,s"S.S.!S.  S.rc  j\VP\4p\VP\4p\ 4'd|V'g V'dl\VP \4'dLV'd\V\ 4'd.M\V\4'd\V\4'gR#VP^,pVP^,pVP^,pVPR ,p VPP^,p \VP P4p WP P8wdR#V ^,^ ,p \WV 4p V fR#VPPV4pVPPW4pVPWg,W4p\ P"!VR R7pRpV V,^, V,pWg,V,W3pV !VVPVP$VP P&.RV 3RV3RV3RV 3RVP$PR ,3.V ^Wg,V,3V ^^3V.\ P(.R 7^,pVPWg,VW4pV^8d\ P*!V^R7pMVP-^4pVPWgW4pVP P/V4p\ P0!V^R7#) zSeparate-kernel decode: fast key scoring + single-tile value weighted sum. 2-3x faster than the fused kernel at large D because each kernel only iterates its own dimensions once.Nrrrrrrrrr)rr/rrrr0rrrr%r0rrr-r7r@rrGrr'rrPrrr)rrPrcrd key_is_prod key_is_mserrrrrrr rr0rrGrjr r rrrrBs&&&& r_separate_score_value_decode.TurboQuantKVCache._separate_score_value_decodes!1EF 0CD    4++-@AA:':;; ,>??<);<<  ! !! $  ! !! $  ! !! $  ! !" %    " "1 %t'',,- '',, ,RB ;HW   >>99/J../?LNN15!/ **YR0 ]*Q.=@ U]*A1 ""$$  ))  " . . 4 4 : :2 >? Q -.Aq $+::,# $ % *%%ae]AA 1 &&3K#++A.K!))!5 !!11+>~~f1--rc V^8dQhRRRR/#rr)rs"rrr,Hs"..!.  .rc \4'd\VP\4'd\VP\4'd\V\ 4'dk\V\ 4'dU\VPP \4'd+\VPP \4'gR#VPpVPpVPpVPpVP PP^8:g&VPPP^8:dR#VP^,pVP^,p VP^,p VP Pp VPPp VP P"P^,p \%WgV 4pVfR#VP&e\(P*!WP&4pV ^,pVRRV 13,P-WW4pVRV V13,P-WW4pVRVVV ,13,P-WW4pVRVV ,R13,P-WW4pEM\(P.!WP0RR7p\(P.!WP2RR7p\(P*!VVP P44p\(P*!VVPP44pVRRV 13,P-WW4pVRV R13,P-WW4pVRRV 13,P-WW4pVRV R13,P-WW4p\7V^, ^4p\7V^, ^4pW,pRpV^,^ ,pV V,^, V,pW,V,V V3pW,V,V ,3p V!VVVVVP P"VP P8VP P:VP P<VP>P"VP>P8VP>P:VP>P<VP P"VP P@VP>P"VP>P@VP PPBVPPPBVP PDVPPDVP PBVPPB.RV 3RV 3RV 3RV3R VP P8PR,3R VP P<PR,3R VP>P8PR,3R VP>P<PR,3R VP P@PR,3RVP>P@PR,3RVP P3RVPP3. ^ VW,V,3RVV V .\(PF\(PF\(PF.R7wp!p"p#V!P-W,VV V4p!V"P-W,VV 4p"V#P-W,VV 4p#\(P6!V#^RR7p$\(PH!V#V$, 4p%\(PJ!V!V%R,,^R7p&\(PJ!V"V%,^R7p'V&\(PL!V'R,\N4, P-WV V4p(\(P*!V(RRV 13,VP PP4p)\(P*!V(RV R13,VPPP4p*\(PR!V)V*.RR7p+\(P.!V+VPTRR7p,\(PV!V,^R7#)usFused decode for SplitCodec — single Metal kernel for score + softmax + weighted_sum across both low/high halves.N.rrDimLowDimHighrrKLMsePackedWidthKLSignPackedWidthKHMsePackedWidthKHSignPackedWidth VLPackedWidth VHPackedWidthVLBitsVHBitsrTrkeepdimsrrr),rrr/rr0rrrrrrrrrrr_rr%rrrr;r@rrrrrFrr^rrrr'rrrOrPrrr8rrr)-rrPrcrdkcvcr rrrrdim_lowdim_highrr rdl2rrrrrrrrrrval_dimr  num_val_tilesr  acc_shapesm_shaperrr global_max scale_factors scaled_accrNrrrrrBs-&&&& r _compiled_split_decode_attention2TurboQuantKVCache._compiled_split_decode_attentionHs   4>>;774++[99:';<<<)=>>4>>335IJJ4++557JKK ^^   ==MM << ! ! & &! +r}}/F/F/K/Kq/P  ! !! $  ! !! $  ! !! $,,""==$$ NN & &q )+HC >  ( ( 4?,I,IJBA+C3=)11!CIC,-55aAGJCsX~!556>>qQQJS#."223;;A!NKGGOZZbAEWW_kkCFYYubll&D&DEFii (G(GHGsHWH}-55aAGIWX .66qQHJ ixi088qKJ!#xy.199!LK8a<+ IM1- $  2", ]*Q.=@ U]*Aw7 EM)A-/$*$$**--((%%++..))  &&  ((!!''!!)) &&// ''00 (( )) %% &&-27#H%" .#Z^^%?%?%E%Eb%IJ$jnn&>&>&D&DR&HI#Z__%@%@%F%Fr%JK$joo&?&?&E&Eb&IJ ,"2"2":":"@"@"DE ,"3"3";";"A"A""EF2<<,,-2==--. mQU]%:;"$h9::rzz2::>U+% !'\//!%7C//!%://!%:VVG!d; w34 VVGmI&>>QG w.Q7!BJJuY/?$FFOO !W ))KXgX 6 8M8MN99[gh79O9OP( 3"=!1!1;~~f1--rc V^8dQhRRRR/#rr)rs"rrr,s"] ] !]  ] rc \4'd\VP\4'dr\VP\ 4'dRVPP P^8d-\V\4'd\V\4'gR#\VPP4pW@PP8wdR#VP^,VP^,VP^,rvpVPR,pVPP^,p V^,^ ,p VPP Pp VPPp V R8:Ed\VVV V RR7p V EeVPPV4wrVP!WVWx4pVP!WVWx4pVPP"pRV3RV3RV 3RVP$PR,3R VP&PR,3R VP(PR,3.pV !VVVPVP$VP*VP&VPVP(VPP P,VPP.VPP,V. V^ ^WV,3RWV,Wx3.\0P2.R 7^,pVP!WVWx4p\0P4!V^R 7#\0P6!WPP84pVR RV13,P!WVWx4pVR VR13,P!WVWx4pRpV V,^, V,p\;WGW4p\=WGV 4pV^,^ ,p^@pVRJ;'d.VV,V,V8;'dV^8;'gV R8pV'EdWV,V,Wx3pWV,V,V,3pV!VVVPVP$VP*VP&VPVP(VPP P,VPP.VPP,. RV3RV3RV3RV 3RVP$PR,3R VP&PR,3R VP(PR,3.^ ^WV,V,3RVVV.\0P2\0P2\0P2.R 7wppp VP!WV,VWx4pVP!WV,VV4pV P!WV,VV4p \0P>!V ^RR7p!\0P@!V V!, 4p"\0PB!VV"R,,^R 7p#\0PB!VV",^R 7p$V#\0PD!V$R,\F4, P!WVWx4p%VPPIV%4p\0P4!V^R 7#VEeV^,^ ,pWV,V,Wx3pWV,V,V,3pV!VVVPVP$VP*VP&VPVP(VPP P,VPP.VPP,. RV3RV3RV3RVP$PR,3R VP&PR,3R VP(PR,3RV3.^ VWV,V,3RVVV.\0P2\0P2\0P2.R 7wppp VP!WV,VWx4pVP!WV,VV4pV P!WV,VV4p \0P>!V ^RR7p!\0P@!V V!, 4p"\0PB!VV"R,,^R 7p#\0PB!VV",^R 7p$V#\0PD!V$R,\F4, P!WVWx4p%VPPIV%4p\0P4!V^R 7#\KV4p&V&!VVPVP$VP*VP&VPVP(VPP8VPP P,VPP.VPP,VPPL4 #)rNrYFrrrrrrrGrr.rirTrValBitsrrr)'rrr/rr0rrrrrr0rr%rrr-r@rrrrr^r'rrrrr;rrrrFrOrPrrrrdr8)'rrPrcrdrrrrrrr rfr ff_kernelr$r]r"rrrBrr r  single_kernel multi_kernelrmin_threadgroups use_singlerrrrrrrrrNrdecodes'&&&& r"_compiled_integer_decode_attention4TurboQuantKVCache._compiled_integer_decode_attentions   4>>+?@@4++-@AA((--1:':;;<);<<4##(() ##(( (  ! !! $  ! !! $  ! !! $  ! !" %    " "1 %RB ~~//44 ""** :2I$ $ > > O  aA1a3 ,,77 AJ"A&"M2& (>(>(D(DR(HI')=)=)C)CB)GH#\%9%9%?%?%CD  "((".."11",,$**$,,009922((11! &a *$%E1=/#%::,'()*Q10~~f155YY(H(H I37 ##A!/CG$$Q10 ]*Q.=@ A ] 4D\J RB   % 4 4!A%)99 4 4!#22qF{  :.5I -13H(5$$**--(( && ((NN,,55NN..$$-- AJ"A&"M2"M2& (>(>(D(DR(HI')=)=)C)CB)GH#\%9%9%?%?%CD!QU]23&((H=!zz2::rzzB5) %GWg<ooae]AAGooae]A>Gooae]A>Ga$?JFF7Z#78M- *B BKJFF7]2;E% 53CT(JJSSaK%%55kBF>>&q1 1  %VNM.5I -13H(4$$**--(( && ((NN,,55NN..$$-- AJ"A&"M2& (>(>(D(DR(HI')=)=)C)CB)GH#\%9%9%?%?%CD%-)>?&((H=!zz2::rzzB5) %GWg<ooae]AAGooae]A>Gooae]A>Ga$?JFF7Z#78M- *B BKJFF7]2;E% 53CT(JJSSaK%%55kBF>>&q1 116      " "  % %      NN , , NN $ $ - - NN & &    % %    % %  rc(V^8dQhRRRRRRRR/#ryr)rs"rrr,s8X,X,X,  X, ! X, X,rc  VeVfVPwr#VPV4pVPV4pVPR,^8wd \R4hVPwrgr\ V\ 4'd(VP PP^,MVPP^,p Wz,p W,PVV V VV 4p VPPp \V4pWP8:EdmVe \ V\4'EdRVR8XEdJ\ VP\4'Ed\ VP\4'Ed\ V\ 4'Ed\ V\ 4'Ed\#VPP$4p\#VPP$4pVP&pVPP)V 4pVPWj,V ,V 4pWg,pVR8:Ed+\+VVV 4pVEeV!VVPVP,VPP.VPVP,VPP..RV 3RV 3RVP,PR,3RVP,PR,3.VR,^^3RVV 3.\0P2.R 7^,pVPWjW4pVPP5V4pVPWgW4P7V4#\9VVV 4p\;4pVEeVEeVR 8:d^@pMVR 8:d^pM VR 8:dR pMRpVV,V 3pVV,3pV!VVPVP,VPP.VPVP,VPP..RV 3RV 3RV3RVP,PR,3RVP,PR,3.V ^ ,Wk,V3^ V ^3VVV.\0P2\0P2\0P2.R 7wppp V!VVV .RV 3RV3.VR,^^3RVV 3.\0P2.R 7^,pVPWjW4pVPP5V4pVPWgW4P7V4#VP=V VV4p!V!e.V!PWgW4pVP7VP&4#VP?V VV4p"V"e.V"PWgW4pVP7VP&4#VPAV VV4p"V"e.V"PWgW4pVP7VP&4#VPP)V 4p#VPPCV#V4p$VPPEV$V4pVPWgW4pVP7VP&4#VPP)V 4p#\0PF!WjWV 3\0P2R7p\0PF!WjW3\0P2R7p%\0PH!WjW3\KR4)\0P2R7p&\M^WP4EF6p'\OVV'VP,4p(\QVV'V(4p)\QVV'V(4p*VPPCV#V)4p$VPSV$V^VV'V(VV4p$VPPUV$V*4wp+p,p-\0PV!V&V-4p.\0PX!V&V., 4p/\0PX!V-V., 4p0VV/R,,V+V0R,,,pV%V/,V,V0,,p%T.p&\0PZ!VV%V&4EK9 V\0PV!V%R,\\4, pVPWgW4pVP7VP&4#)Nz9TurboQuant decode attention expects a single query token.rnr)rrrFrGrri rYir*iBlocksrr{r?r)rrnrnr)/rr`rrrrrr%r@r0r_rr|rr/rrr0rrr-r4rr'rrrrr:r@rrrr7rHrrrrOrrrvrOrrOrWr)1rr-rcrdr`rrrrrrrrPrrlrrrr$ q_rot_flatBQH fused_kernelrrrBpass1pass2 num_blocksrrrr8r9 sep_outputrBr0rrrrirjrrrrrrrrs1&&&&&& rdecode_attention"TurboQuantKVCache.decode_attentions  !5'+zz $J\\*- ||L1 ==  !K %]]a*&:;; NN & &q )!!''*  + "?33    $$(( $Z0 55 5 LZc22tx7G4>>+>??t//1DEEz+=>>|-?@@t~~223t//445 66G"]]1>I+EqI m4'#;HhPQ#RL#/* * * 0 0 * 2 2 $ 7 7 , 2 2 , 4 4 $ 0 0 9 9$"' !. :!/1C1C1I1I"1M N!/1E1E1K1KB1O P & #&*a!3(4,/8*+-::,'()*'*kk!&N !%!1!1!A!A+!N%~~aAIPPQVWW98QO8:$):#t+%' %.%( %.%( %( !$z!11 5I #j 02H27&&,,&.. NN33(..(00 ,,55 #AJ*I6%z2+Z-?-?-E-Eb-IJ+\-A-A-G-G-KL "'OM& &(A$6'0(H&E')zz2::rzz&J13/GXx6 '8<#(!*x.D!E!Dj!Q/$0(+Qxj')zzl C#&++aY"JK!--==kJF!>>!ELLUSS::J %#++A!G}}W]]33AAK &$,,Q1H}}W]]33??K &$,,Q1H}}W]]33#~~==oN ^^223CZPF%%>>v|TF^^A!?F==/ />>99/J1) B"**UXXqi;2::N GG I ) 5\M** Q .H.HIG g0J0J&JKE*:wFI,\7EJK^^223CYOF// F  ?? T 1L+yjjI6G G 34J&&W!45KI.. I@V1VV $j0;3LLJI GGFJ 2;J>"**Z %:DAAa;}}W]]++rc VP#rrr"s&rr/TurboQuantKVCache.sizes {{rc "VPfR#VPVP8Xd VP#\ VPVP4\ VP VP43pWnVPVnV#)NrH)r.r2rr1rrV)rsliceds& rrTurboQuantKVCache.statesv 99    $ $ 3%% %dii5| KK8  $$(KK! rc RVnRVnVfRRuVnVn^VnR#VwVnVn\ VP4VnR#rU)r1r2r.rVrrrrs&&rrrsO!$&! =%)4 "DIt{DK !& 4;#DII. rc v\\\VPVPVP 344#r)rmaprrrrr"s&r meta_stateTurboQuantKVCache.meta_states'St{{DIItyyABCCrc \V^,4Vn\V^,4Vn\V^,4VnR#rN)r0rrrrrs&&rrrs/%(m %(O aM rc R#)Trr"s&r is_trimmableTurboQuantKVCache.is_trimmablesrc \VPV4pV;PV,unRVnRVnV#rU)rrr1r2rr%s&&rtrimTurboQuantKVCache.trims5  Q  q !$&!rc 0\VRVP/VB#r)rr)rargskwargss&*,r make_maskTurboQuantKVCache.make_masks$dI4;;I&IIrc VPRJ#rr.r"s&rr_TurboQuantKVCache.emptyyyD  rc ,\VP4#r)rrr"s&rrTurboQuantKVCache.nbytessTZZ((r) r1r2r3r4rr/r.rrr0rVi@rH)NNr.N)$rrrrr|r}r~rVDEFAULT_TURBOQUANT_SEEDr classmethodr=rCrQr9 staticmethodr`r)rvrrrrrrr/propertyrsetterrrrrr_rrrrrr(r(Ps#!!J0G #-D*7 r% NRR#>],~DHLS.j.B] ~X,t   \\//DD"" J!))rr(c]tRtRtRtRt]3RRlltRRltRR lt R R lt R R lt RRlt ] R4t]PR4t] R4t]PR4tRtRtRtRRRllt] R4t] R4tRtR#)BatchTurboQuantKVCacheiafBatch-aware TurboQuant KV cache for continuous batching. Wraps TurboQuant's quantization codecs with per-sequence offsets and left-padding so that ``Batch.extend`` / ``Batch.filter`` work during continuous batching. Unlike ``BatchQuantizedKVCache`` (uniform ``mx.quantize``), this uses TurboQuant's MSE/Prod codecs for higher quality at the same bit-rate. r*c$V^8dQhRRRRRR/#)r left_paddingrrrrr0r)rs"rr#BatchTurboQuantKVCache.__annotate__s( rc \V4VnW0nRVnRVnRVnRVn\P!V4Vn \P!VUu.uFqD)NK up4Vn ^Vn R#uupir) rrrr.rVr/r0rrrr_idx)rrrrlps&&&& rrBatchTurboQuantKVCache.__init__sn #4(    HH\2hhl;ll;<   > %  JJrNTYYTYY? 4&??rc V^8dQhRRRR/#r@r)rs"rrrs" " X" x" rc <VPV4VPpVPPV4pVPPV4pW1P ^,,pVP f"\WF4Vn\WV4VnML\VP W6VP4Vn\VPW6VP4Vn\VP WC4\VPWS4V;PVP ^,, un W`nVP ^,^8gVP^2,^8Xd,\P!VP VP4\VP VP4p\VPVP4pVP ^,p \!WpPV 4\!WPV 43#r )rCr r/r%r0rr.rrVrrVrrrrWrr) rr.rVprevrXrYrZr\r]rs &&& rr9'BatchTurboQuantKVCache.update_and_fetchsv D!yy>>**40%%..v6 A& 99 ,X?DI.zCDK/ 4$//DI2 TDOODK TYY/T[[*3 tzz!}$  ::a=1 R1!4 GGDIIt{{ + $))TYY / $++tyy 1**Q- YY 8 YY 8  rcV^8dQhRR/#rr)rs"rrrs((H(rc JaVPe7\VPV4Vn\VPV4VnVPV,VnVPV,VnVPP 4P 4oS^8dVPe=V3Rlp\VPV4Vn\VPV4VnV;PS,unV;PS,unR#R#)NcH<V^8XdVRSR13,#VRSR1R3,#)r.Nr~r)rrmin_lps&&r_trim,BatchTurboQuantKVCache.filter.._trim!s-qy fg.S&'1_--r) r.rrVrrritemrr )rrrrs&& @rfilterBatchTurboQuantKVCache.filters 99 %dii?DI' ]CDKkk-0  --m<""&&(--/ A:yy$. 'tyy%8 (e< II I    '  rcV^8dQhRR/#)r otherz'BatchTurboQuantKVCache'r)rs"rrr+s1141rc a VPfsVPfe\P!VPVP.4Vn\P!VPVP.4VnR#\ VP VP 4o V 3RlpV!V4pV!V4pVfiVfe\P!VPVP.4Vn\P!VPVP.4VnR#VfoVwVnVnrV\P!VPV.4Vn\P!VPS ,V.4VnS VnR#VfmVwVnVnrV\P!WQP.4Vn\P!WaPS ,.4VnS VnR#VwrxrVVwrr\Wy4Vn\W4Vn\P!W[.4Vn\P!Wl.4VnS VnR#)Nc:<VPfR#SVP, p^p\\VPVP4W4p\\VPVP4W4pVP V,pW4VP V3#r)r.r rrrVrr) cache_objrrkrr max_idxs& r _pad_side0BatchTurboQuantKVCache.extend.._pad_side3s~~%Y^^+DE!Y^^Y^^A''$.B))2- -r) r.rrrrrFr rVr)rrr$r_selfr_othersoslpsksvokovooolpr#s&& @rrBatchTurboQuantKVCache.extend+s 99 !3 "0A0A5CUCU/V WD ..$++u||)DEDK dii, .4E" >go "0A0A5CUCU/V WD ..$++u||)DEDK  >.5 +DIt{B..$++r):;DK "0A0AG0KS/Q RD DI  ?.4 +DIt{B.."ll); 4; T-> 99 %dii0DI !rc v\\\VPVPVP 344#r)rrrr rrr"s&rr!BatchTurboQuantKVCache.meta_states'Styy$))TYY?@AArc \V^,4Vn\V^,4Vn\V^,4VnR#r)r0r rrr)rrs&&rrr8s/!I !A$K !I rc R#Frr"s&rr#BatchTurboQuantKVCache.is_trimmablesrc ^#rrrs&&rrBatchTurboQuantKVCache.trimsrc VPRJ#rrr"s&rr_BatchTurboQuantKVCache.emptyrrc V^8dQhRRRR/#)r Nr0 return_arrayrr)rs"rrrs  3 d rc J\V3RVPRVP/VB#)rr)r r r)rrBrCrs&&&,rr BatchTurboQuantKVCache.make_masks4!  ii .2.?.? CI  rc ^@#)@rr"s&r group_size!BatchTurboQuantKVCache.group_sizesrc VPf^#\VPVP4p\VPVP4p\ V4\ V4,#r)r.rr rVr)rsrs& rrBatchTurboQuantKVCache.nbytessL 99  DII . dii 0Q-"222r) r rr/r.rrrr0rVrHr;)rrrrr%rVrrrCr9rrr)rrrrrrr_rrHrrrrrrrsJ , (+" P(*1n  44 \\11 BB !  33rrr)r)rF)r2)r*)rr;)n __future__rr functoolsrtypingrrmlx.corecorernumpyrmlx_lm.models.cacherrr rrrrr,rSrWrZrarfrkrrrrrrrrrr0r5r<rCrHrKrdrrrrrrrrr4r:r@r]rroryrrrrrrrrrrrrrr r#r(r+r2r@r:rErVrirqrvrzrrrrrrrrrrrrrrrrrRrrrrrrr(rrrrrTsT" 'UU < 411h 4[[| 4((V 4> 400f 4HHV 444n 4kk\ 4JJZ 4~~B 4GGT 4TTn 4PPf8*v*DI*X+*\8*v-"G*T6*r5*p 4//d 4MM` 4AAH 4==@ 4UUp 4NNb4n. 4CCL 4ddN 4PPf!HD 4yyx 4IIXXv 4TTn 4QQh 4ppf *(:( z: =47 & 4 * *5 4## 4    4//$ 4$2$2N  4&2&2R& $>N$:N%IPK@ KF   K 9KxI@$KN HF? #R k)k)\_D_DDtItInBIBIJ*aOaOH1!!*i) i)^#e3Ze3r