+ TV-jg^RIHt^RIHtRtRtRt]!4t]!4t ]PR4t ] PR4t Rt RtR t]!4t]!4t]PR 4t]PR 4tR tR#) Nc\P!4\P8H;'d\PP 4#)N)mxdefault_devicegpumetal is_availabled/Volumes/fast/ai/experiments/ui-tars-smoke/.venv/lib/python3.14/site-packages/mlx_lm/tuner/losses.py can_run_metalr s/   "&& ( D DRXX-B-B-DDr cx\4'gR#Rp\PPRRR.R.VRR7#)Naz constexpr int M = 4; constexpr int block = 1024 * M; constexpr int full_blocks = V / block; constexpr int extra = V - full_blocks * block; threadgroup float shared[32 * 2]; uint out_idx = threadgroup_position_in_grid.y; uint simd_lane_id = thread_index_in_simdgroup; uint simd_group_id = simdgroup_index_in_threadgroup; logits_q += out_idx * V; logits_p += out_idx * V; out += out_idx; float lse_q_minus_p; float lse_p; { float max_q = -1e30; float max_p = -1e30; float sum_exp_q = 0; float sum_exp_p = 0; int offset = thread_index_in_threadgroup * M; for (int i = 0; i < full_blocks; i++) { // Read and update q and p float vals_q[M]; float vals_p[M]; for (int j=0; j 0) { // Read and update q and p float vals_q[M]; float vals_p[M]; for (int j=0; j < M; j++) { vals_q[j] = (offset + j < V) ? logits_q[offset + j] : -1e30; vals_p[j] = (offset + j < V) ? logits_p[offset + j] : -1e30; } float prev_max_q = max_q; float prev_max_p = max_p; for (int j=0; j 0) { float vals_q[M]; float vals_p[M]; for (int j=0; j(kl); } } kl_forwardlogits_qlogits_poutTname input_names output_namessourceensure_row_contiguousr rfast metal_kernelrs r _make_kl_forward_kernelr sI ??ZFx 77   ,W"  r cx\4'gR#Rp\PPR.ROR.VRR7#)Na constexpr int M = 4; constexpr int block = 1024 * M; constexpr int full_blocks = V / block; constexpr int extra = V - full_blocks * block; threadgroup float shared[32 * 2]; uint out_idx = threadgroup_position_in_grid.y; uint simd_lane_id = thread_index_in_simdgroup; uint simd_group_id = simdgroup_index_in_threadgroup; logits_q += out_idx * V; logits_p += out_idx * V; out += out_idx * V; cotan += out_idx; float lse_q; float lse_p; { float max_q = -1e30; float max_p = -1e30; float sum_exp_q = 0; float sum_exp_p = 0; int offset = thread_index_in_threadgroup * M; for (int i = 0; i < full_blocks; i++) { // Read and update q and p float vals_q[M]; float vals_p[M]; for (int j=0; j 0) { // Read and update q and p float vals_q[M]; float vals_p[M]; for (int j=0; j < M; j++) { vals_q[j] = (offset + j < V) ? logits_q[offset + j] : -1e30; vals_p[j] = (offset + j < V) ? logits_p[offset + j] : -1e30; } float prev_max_q = max_q; float prev_max_p = max_p; for (int j=0; j( c * (metal::fast::exp(vals_q[j] - lse_q) - metal::fast::exp(vals_p[j] - lse_p))); } // Move to the next block offset += block; } if (extra > 0) { float vals_q[M]; float vals_p[M]; for (int j=0; j( c * (metal::fast::exp(vals_q[j] - lse_q) - metal::fast::exp(vals_p[j] - lse_p))); } } } } kl_backwardrTr)rrcotanrrs r _make_kl_backward_kernelr sD ??SFj 77   5W"  r c VPVPR,,pVPp\W.VPRR.V.RV3RVPR,3.RV^3RR7^,#NTVinputs output_shapes output_dtypestemplategrid threadgroupr&r#r#)sizeshapedtype_kl_forward_kernel)rrn_outsdts&& r _kl_div_lossr6Xs| ]]hnnR0 0F B #~~cr*+d)c8>>"#567FA     r c Vwr4VPp\P!V4p\W4V.VP.V.RV3RVPR,3.RVP ^3RR7^,pWv3#r$r%r&r'r.r/)r2r zeros_like_kl_backward_kernelr1r0)primals cotangentoutputrrr5dpdqs&&& r r6r6gs H B x B I.~~&d)c8>>"#567INNA &   B 6Mr c \4'd \W4#\PP V\ P !VRRR7, V\ P !VRRR7, RRR7#)r#TaxiskeepdimsnonerB reductionr.)r r6nnlosses kl_div_lossr logsumexp)rrs&&r rIrIys`H//yy$$ r||H2E E r||H2E E %  r cz\4'gR#Rp\PPRRR.RR.VRR7#) NaT constexpr int M = 4; constexpr int block = 1024 * M; constexpr int full_blocks = V / block; constexpr int extra = V - full_blocks * block; threadgroup float shared[32 * 2]; uint out_idx = threadgroup_position_in_grid.y; uint simd_lane_id = thread_index_in_simdgroup; uint simd_group_id = simdgroup_index_in_threadgroup; logits_q += out_idx * V; logits_p += out_idx * V; out += out_idx; out_kl_q += out_idx; float lse_p; float lse_q; { float max_q = -1e30; float max_p = -1e30; float sum_exp_q = 0; float sum_exp_p = 0; int offset = thread_index_in_threadgroup * M; for (int i = 0; i < full_blocks; i++) { // Read and update q and p float vals_q[M]; float vals_p[M]; for (int j=0; j 0) { // Read and update q and p float vals_q[M]; float vals_p[M]; for (int j=0; j < M; j++) { vals_q[j] = (offset + j < V) ? logits_q[offset + j] : -1e30; vals_p[j] = (offset + j < V) ? logits_p[offset + j] : -1e30; } float prev_max_q = max_q; float prev_max_p = max_p; for (int j=0; j(2)); int offset = thread_index_in_threadgroup * M; for (int i = 0; i < full_blocks; i++) { // Read and add to the kl_p and kl_q float vals_q[M]; float vals_p[M]; for (int j=0; j 0) { float vals_q[M]; float vals_p[M]; for (int j=0; j(0.5 * kl_p + 0.5 * kl_q); out_kl_q[0] = static_cast(kl_q); } } js_forwardrrrout_kl_qTrrrs r _make_js_forward_kernelrNsL ??lF\ 77   ,Z("  r cx\4'gR#Rp\PPR.ROR.VRR7#)Na constexpr int M = 4; constexpr int block = 1024 * M; constexpr int full_blocks = V / block; constexpr int extra = V - full_blocks * block; threadgroup float shared[32 * 2]; uint out_idx = threadgroup_position_in_grid.y; uint simd_lane_id = thread_index_in_simdgroup; uint simd_group_id = simdgroup_index_in_threadgroup; logits_q += out_idx * V; logits_p += out_idx * V; out_q += out_idx * V; cotan += out_idx; output_kl_q += out_idx; float lse_q; float lse_p; { float max_q = -1e30; float max_p = -1e30; float sum_exp_q = 0; float sum_exp_p = 0; int offset = thread_index_in_threadgroup * M; for (int i = 0; i < full_blocks; i++) { // Read and update q and p float vals_q[M]; float vals_p[M]; for (int j=0; j 0) { // Read and update q and p float vals_q[M]; float vals_p[M]; for (int j=0; j < M; j++) { vals_q[j] = (offset + j < V) ? logits_q[offset + j] : -1e30; vals_p[j] = (offset + j < V) ? logits_p[offset + j] : -1e30; } float prev_max_q = max_q; float prev_max_p = max_p; for (int j=0; j(2)); float kl_q = output_kl_q[0]; int offset = thread_index_in_threadgroup * M; for (int i = 0; i < full_blocks; i++) { // Read and compute vjp for logits_q float vals_q[M]; float vals_p[M]; for (int j=0; j( c * 0.5 * q_j * (logtwo - metal::fast::log(1 + metal::fast::exp(logp_j - logq_j)) - kl_q) ); } // Move to the next block offset += block; } if (extra > 0) { float vals_q[M]; float vals_p[M]; for (int j=0; j( c * 0.5 * q_j * (logtwo - metal::fast::log(1 + metal::fast::exp(logp_j - logq_j)) - kl_q) ); } } } } js_backwardout_qTr)rrr output_kl_qrrs r _make_js_backward_kernelrS?sD ??]F~ 77   DY"  r c @VPVPR,,pVPp\W.VPRRVPRR.W3.RV3RVPR,3.RV^3RR7pV^,\P !V^,43#r")r0r1r2_js_forward_kernelr stop_gradient)rrr4r5outputss&& r _js_div_lossrXs ]]hnnR0 0F B #~~cr*HNN3B,?@h)c8>>"#567FA  G 1:r'' 3 33r c Vwr4VwrVVwrgVPp\P!V4p \W4WW.VP.V.RV3RVPR,3.RVP ^3RR7p W3#r8)r2rr9_js_backward_kernelr1r0) r; cotangentsrWrrr_kl_qr5r>r?s &&& r rXrXs HHEGA B x B E0~~&d)c8>>"#567EJJ "  B 6Mr cP\4'd\W4^,#V\P!VRRR7, pV\P!VRRR7, pV\P!^\P !W2, 4,4,\P!^4P VP4, p\PPWBRRR7p\PPWCRRR7pRWV,,#)rTrArDrEg?r.) r rXrrJlogexpastyper2rGrHrI)rr logprobs_p logprobs_q logprobs_mkl_pr]s&& r js_div_lossrfsH/22 XB NN  XB NN ffQ 7889 :ffQix~~. /  yy$$Z"PV$Wyy$$Z"PV$Wdk""r )mlx.corecorermlx.nnrGr rr r3r:custom_functionr6vjprIrNrSrUrZrXrfr r r rlsEeP^B-..0    "  wthV-..0 4 4$ #r