+ UV-j+P^RIHt]PP 4tRtR RltR RltR Rlt R Rlt Rt Rt R t R#)NcZVPwr4rV\P!\P!V\PR7WQ, ,4P \P 4p\P!\P!V\PR7Wb, ,4P \P 4p\P!V^V^, 4p\P!V^V^, 4pVRRV3,RRRV3,#)z:Pure-MLX nearest-neighbor interpolation (no Metal kernel).dtypeNNN)shapemxfloorarangefloat32astypeint32clip) xout_hout_w batch_sizechannelsin_hin_wy_idxx_idxs &&& g/Volumes/fast/ai/experiments/ui-tars-smoke/.venv/lib/python3.14/site-packages/mlx_vlm/models/kernels.py_nearest_interpolate_mlxrs'(ww$J$ HHRYYuBJJ74<H I P P  E HHRYYuBJJ74<H I P P  E GGE1dQh 'E GGE1dQh 'E Q5[>!Q5. ))c^VPwr4rVVeVwrxMUVeG\V\\34'dT;rMVwr\WY,4\Wj,4rM \ R4h\ 'g \ WV4#\P!W4WVWx.\PR7p VPR 4p VPp V \P8wd V P\P4p Rp\PPRRR.R.VR7p\!W4pV!W.WW4,3VW4,V,V,3.\P.R 7pV^,PW4Wx4pV \P8wdVPV 4pV#) zI Nearest neighbor interpolation that exactly matches PyTorch's behavior. -Either size or scale_factor must be specifiedra uint x_out = thread_position_in_grid.x; uint y_out = thread_position_in_grid.y; uint bc_idx = thread_position_in_grid.z; int batch_size = dims[0]; int channels = dims[1]; int in_h = dims[2]; int in_w = dims[3]; int out_h = dims[4]; int out_w = dims[5]; if (x_out >= (uint)out_w || y_out >= (uint)out_h || bc_idx >= (uint)(batch_size * channels)) return; int c = bc_idx % channels; int b = bc_idx / channels; // PyTorch's coordinate calculation for nearest neighbor // This matches: torch.nn.functional.interpolate(..., mode='nearest') float scale_h = float(in_h) / float(out_h); float scale_w = float(in_w) / float(out_w); // PyTorch uses floor for nearest neighbor coordinate mapping int y_in = int(floor(float(y_out) * scale_h)); int x_in = int(floor(float(x_out) * scale_w)); // Clamp to bounds y_in = max(0, min(y_in, in_h - 1)); x_in = max(0, min(x_in, in_w - 1)); int input_offset = ((b * channels + c) * in_h + y_in) * in_w + x_in; int output_offset = ((b * channels + c) * out_h + y_out) * out_w + x_out; output[output_offset] = input[input_offset]; nearest_interpolationinputdimsoutputname input_names output_namessourceinputsgrid threadgroup output_shapes output_dtypes)r isinstanceintfloat ValueError _HAS_METALrrarrayr reshaperr r fast metal_kernelget_optimal_threadgroup)rsize scale_factorrrrrrrscale_hscale_wrx_flat input_dtyper%kernelr)outputsresults&&& rnearest_interpolater@s()ww$J$  u  ! lS%L 1 1 , ,Gg+ G4>*C,?uHII :'%88 88Z4uDBHH UDYYr]F''Kbjj rzz*#FJWW ! ! $f%Z "F*%7K~J1 2"-5=?@zzl GQZ   e CFbjj {+ Mrc \P!V4pW",pW2,pVR,V,VR,V,, R,pW,RV,V,, RV,V,,RV,, p\P!VR8*V\P!VR8V\P!V444#)a7Cubic convolution kernel. PyTorch's ``F.interpolate(mode="bicubic")`` uses ``a=-0.75`` (ATen ``upsample_bicubic2d``) for the non-antialiased path, but its antialiased path matches Pillow, which uses ``a=-0.5`` (Keys' cubic). The caller selects the coefficient per axis based on whether antialiasing is active. @g@?g@g @g@)rabswhere zeros_like)taatat2at3w1w2s&& r _cubic_weightrNqs B 'C (C c'SAGs? *S 0B 37S= 37R< /#' 9B 88B#Ir288BHb"--:J#K LLrcaVPwrVrxVPp VP\P4pW, p W(, p V'dV^8dV^8d\P !V\PR7V^, ,V^, , p \P !V\PR7V^, ,V^, , p M\P !V\PR7R,V, V,R, p \P !V\PR7R,V, V,R, p S'dV R8d RV , MRpS'dV R8d RV , MRpRV,pRV,pV3RlpV!WVV4wppV!WVV4wppVRRVP R4R3,pVP^,pVP^,pVP WVVVV4p\P!VVR ,,^R7pVRRRVP R43,pVP WVWV4p\P!VVR ,,^R7pV \P8wdVPV 4pV#) z1Pure-MLX bicubic interpolation (no Metal kernel).rg?rCrBc<\P!W, 4P\P4^,p\ ^V,^,4p\P !V\PR7pVR,VR,,pVR,VP\P 4, p\W, S 'dRMR4p V^8Wq8,p W,p \P!V^V^, 4pV \P!V R RR7R,, p Wy3#) rT)axiskeepdimsg:0yE>rN)Nrgr,) rr r r r.r r rNrsum) coordsin_sizefssupportstartn_tapsoffsetspixdistwmask antialiass &&&& r _weights_1d-_bicubic_interpolate_mlx.._weights_1ds)*11"((;a?Q[1_%))F"((3Gnww//gBJJ!77 $)YTE BqS]+ Hggc1gk* T2T9 :v rr)rRr,)NNrrN)NNNrr)rrr rr r r3rV)rrr align_cornersrbBCrrr<r9r:y_outx_outfs_hfs_w support_h support_wrcpix_ywypix_xwx gathered_ythtwtmp gathered_xr?s&&&&f r_bicubic_interpolate_mlxrws!wwA$''K AlGlGuqy %rzz2dQh?519M %rzz2dQh?519M5 3c9UBTICO5 3c9UBTICO )Ws]C'MD(Ws]C'MDd Id I"Ey9IE2Ey9IE21ar*A-.J QB QB##A%T:J &&b!788q ACQ1emmB//0J##A%;J VVJ$:!;;! DFbjj {+ Mrc ~VPwrVrxVeVwrW, W, rMUVeG\V\\34'dT;rMVwr\W{,4\W,4rM \ R4h\ 'g\ W WV4#Rp V'dRMRpV'dV R8d RV , pMRpV'dV R8d RV , pMRp\P!T T V'dRMRVVVV .\PR7p\P!WVWxW.\PR7pVPR 4pVPpV\P8wd VP\P4pRpRp\PPR.R OR .VVR 7p\!W4pV!VVV.WWV,3VWV,V ,V ,3.\P.R 7pV^,PWVW4pV\P8wdVPV4pV#)aY Bicubic interpolation using MLX's built-in interpolate function. Args: x: MLX tensor of shape [B, C, H, W] size: Tuple of (out_h, out_w) or None scale_factor: Float or tuple of (scale_h, scale_w) or None align_corners: Whether to align corners antialias: Whether to apply antialiasing Returns: Interpolated MLX tensor rrBrCgra // Cubic convolution kernel. PyTorch's non-antialiased bicubic uses // a=-0.75 (ATen upsample_bicubic2d); its antialiased path matches // Pillow, which uses a=-0.5 (Keys' cubic). float cubic_kernel_a(float x, float a) { float absx = fabs(x); float absx2 = absx * absx; float absx3 = absx2 * absx; if (absx <= 1.0f) { return (a + 2.0f) * absx3 - (a + 3.0f) * absx2 + 1.0f; } else if (absx < 2.0f) { return a * absx3 - 5.0f * a * absx2 + 8.0f * a * absx - 4.0f * a; } return 0.0f; } float cubic_kernel(float x) { return cubic_kernel_a(x, -0.75f); } // Antialiased bicubic kernel - widens the support for downsampling and // uses Pillow's a=-0.5 to match PyTorch's antialiased path. float cubic_kernel_antialias(float x, float scale) { return cubic_kernel_a(x / scale, -0.5f); } aH // Get thread position uint x_out = thread_position_in_grid.x; uint y_out = thread_position_in_grid.y; uint bc_idx = thread_position_in_grid.z; // Extract dimensions int batch_size = dims[0]; int channels = dims[1]; int in_h = dims[2]; int in_w = dims[3]; int out_h = dims[4]; int out_w = dims[5]; // Extract parameters float scale_h = params[0]; float scale_w = params[1]; bool align_corners = params[2] > 0.5f; bool use_antialias = params[3] > 0.5f; float filter_scale_h = params[4]; float filter_scale_w = params[5]; float support = params[6]; // Check bounds if (x_out >= (uint)out_w || y_out >= (uint)out_h || bc_idx >= (uint)(batch_size * channels)) return; // Calculate batch and channel indices int c = bc_idx % channels; int b = bc_idx / channels; // Calculate input coordinates float x_in, y_in; if (align_corners && out_w > 1 && out_h > 1) { x_in = float(x_out) * (in_w - 1) / (out_w - 1); y_in = float(y_out) * (in_h - 1) / (out_h - 1); } else { // PyTorch's default coordinate mapping x_in = ((float(x_out) + 0.5f) / float(out_w)) * float(in_w) - 0.5f; y_in = ((float(y_out) + 0.5f) / float(out_h)) * float(in_h) - 0.5f; } // Calculate the support region based on antialiasing float support_h = use_antialias ? support * filter_scale_h : support; float support_w = use_antialias ? support * filter_scale_w : support; // Calculate the range of input pixels to sample int y_start = int(floor(y_in - support_h)) + 1; int y_end = int(floor(y_in + support_h)) + 1; int x_start = int(floor(x_in - support_w)) + 1; int x_end = int(floor(x_in + support_w)) + 1; // Clamp to valid range y_start = max(0, y_start); y_end = min(in_h, y_end); x_start = max(0, x_start); x_end = min(in_w, x_end); // Perform bicubic interpolation with antialiasing float result = 0.0f; float weight_sum = 0.0f; for (int y_pos = y_start; y_pos < y_end; y_pos++) { float dy = float(y_pos) - y_in; float wy = use_antialias ? cubic_kernel_antialias(dy, filter_scale_h) : cubic_kernel(dy); for (int x_pos = x_start; x_pos < x_end; x_pos++) { float dx = float(x_pos) - x_in; float wx = use_antialias ? cubic_kernel_antialias(dx, filter_scale_w) : cubic_kernel(dx); float weight = wy * wx; // Calculate input tensor offset int input_offset = ((b * channels + c) * in_h + y_pos) * in_w + x_pos; // Add weighted contribution result += input[input_offset] * weight; weight_sum += weight; } } // Normalize by weight sum if (weight_sum > 1e-8f) { result /= weight_sum; } // Calculate output tensor offset int output_offset = ((b * channels + c) * out_h + y_out) * out_w + x_out; // Assign the result to output output[output_offset] = result; bicubic_interpolation_antialiasr )r"r#r$r%headerr&r,)rrparams)rr-r.r/r0r1rwrr2r r r3rr r4r5r6)rr7r8rerbrrrrrrr9r:rZantialias_flagfilter_scale_hfilter_scale_wr{rr;r<rzr%r=r)r>r?s&&&&& rbicubic_interpolaters"()ww$J$   <  ! lS%L 1 1 , ,Gg+ G4>*C,?uHII :'% RRG &S3NWs]wWs]wXX   Cc     jj F 88Z4uDBHH UDYYr]F''Kbjj rzz*F8`FFWW ! ! ./Z "F*%7Kf%J1 2"-5=?@zzl GQZ   e CFbjj {+ Mrc:aaaaaaaSPwooooVPwpoopVR,^,S,^, ^, pVR,^,S,^, ^, p\P!V4P\P4p\P!V4P\P4pV^,pV^,pVP\P 4V, VP\P 4V, ,R,p W5P\P 4, VP\P 4V, ,R,p VP\P 4V, WFP\P 4, ,R,p W5P\P 4, WFP\P 4, ,R,p VVVVVVV3Rlp W!We4,W!Wg4,,W!W4,,W!W4,,pV#)uPure-MLX bilinear grid sample (no Metal kernel). x: (B, H, W, C) — channel-last grid: (B, gN, gM, 2) — normalised [-1, 1] c<V^8VS 8,V^8,VS 8,p\P!V^S ^, 4p\P!V^S ^, 4pVS ,V,PS R4pSPS S S ,S 4p\P!S 4R,pWgV3,pVPS SS S 4pWR,,#)rr,rT.N)rrr3r )yyxxvalidyy_cxx_cidxr;b_idxvalsrfrgHWgMgNrs&& r_gather!_grid_sample_mlx.._gathersqR!V$a0BF;wwr1a!e$wwr1a!e$ax$''2.1a!eQ' ! W%Sj!||Ar2q)I&&&r).r).rQr)rrr r r r )rr(_ixiyix0iy0ix1iy1wawbwcwdroutrfrgrrrrsf& @@@@@@r_grid_sample_mlxrs JAq!Q::LAr2q = 0 && iy_nw <= H - 1 && ix_nw >= 0 && ix_nw <= W - 1 ? I_nw : 0; I_ne = iy_ne >= 0 && iy_ne <= H - 1 && ix_ne >= 0 && ix_ne <= W - 1 ? I_ne : 0; I_sw = iy_sw >= 0 && iy_sw <= H - 1 && ix_sw >= 0 && ix_sw <= W - 1 ? I_sw : 0; I_se = iy_se >= 0 && iy_se <= H - 1 && ix_se >= 0 && ix_se <= W - 1 ? I_se : 0; out[elem] = nw * I_nw + ne * I_ne + sw * I_sw + se * I_se; grid_samplerr(rr!T)r'templater*r+r(r))rQrQ) ndimrr1rrr4r5rprodr2) rr(rfrrgrrD out_shaper%r=r>s && rrrs 66Q;))); 99>///>JAq!::LA2I 67776 :((0FdWW ! ! &MW "Fy.! kwwiggbhhy)*Aq 1 G 1:rcRpRpRp\W04p\W14p^VP4^, ,pWu8d V^,p^VP4^, ,pW8d V^,pWx,V8dWx8d V^,pKV^,pK)\^V4p\^V4pWx^3# \dTu#i;i)i) rrQ)min bit_lengthmax Exception) rrMAX_THREADS_PER_GROUPMAX_THREADS_PER_DIMdefault_threadgroup max_width max_heightwidthheights && rr6r62s !&#+3 ,4 i**,q01  QJEz,,.23  q[Fn44 1Au Qq!! #""#sB4B== C  C )NN)rU)FF)NNFF)mlx.corecorermetal is_availabler1rr@rNrwrrrr6rrrsK XX " " $ *"Wt M <~fR+\Xv)#r