+
    KV-jˆ	  ã                   ó�   € ^ RI HtHtHt ^ RIt^ RIHt ^RI	H
t
 ^RIHtHtHt ^ RIHt ]P                   R R l4       tR R	 ltR# )
é    )ÚTupleÚListÚcastN)Ú	LLMatcher)Úget_bitmask_shapeÚallocate_token_bitmaskÚfill_next_token_bitmask)ÚNDArrayc                óx   € V ^8„  d   QhR\         P                  R\         P                  R\         P                  /# )é   ÚdataÚmaskÚreturn)ÚmxÚarray)Úformats   "Ú_/Volumes/fast/ai/experiments/ui-tars-smoke/.venv/lib/python3.14/site-packages/llguidance/mlx.pyÚ__annotate__r      s-   € ÷ 'ñ '¤R§X¡Xð '´R·X±Xð 'Ä"Ç(Á(ñ 'ó    c           	      ór  € R p\         P                  P                  R. ROR.VR7      p\         P                  ! \	        R4      ) .V P
                  R7      pV! WV.RV P
                  3.V P                  ^,          V P                  ^ ,          ^3R	V P                  .V P
                  .R7      pV^ ,          pV# )
a¹  
        uint batch = thread_position_in_grid.y;  // Batch index
        uint elem = thread_position_in_grid.x;   // Element index within batch

        // Bounds check to prevent out-of-bounds access
        // assert(batch < inp_shape[0] && elem < inp_shape[1]);

        uint word_idx = elem / 32;  // Which u32 word
        uint bit_idx = elem % 32;   // Which bit in the word

        // Bounds check for mask access
        // assert(word_idx < mask_shape[1] && batch < mask_shape[0]);

        uint bit = word_idx < mask_shape[1] && (mask[batch * mask_shape[1] + word_idx] >> bit_idx) & 1;
        out[batch * inp_shape[1] + elem] = bit ? inp[batch * inp_shape[1] + elem] : neg_inf[0];
    Úbitmask_apply_batchedÚout)ÚnameÚinput_namesÚoutput_namesÚsourceÚinf)ÚdtypeÚT)ÚinputsÚtemplateÚgridÚthreadgroupÚoutput_shapesÚoutput_dtypes)Úinpr   Úneg_inf)é   é   r)   )r   ÚfastÚmetal_kernelr   Úfloatr   Úshape)r   r   r   Úkernelr'   ÚoutputsÚas   &&     r   Úapply_token_bitmask_kernelr1      s³   € ð€Fô" �W‰W×!Ñ!Ø$Ú.Ø�WØð	 "ó €Fô �hŠhœ˜u›˜�¨d¯j©jÔ9€GáØ˜GÐ$Ø˜Ÿ
™
Ð#Ð$Ø�j‰j˜�m˜TŸZ™Z¨�]ØðàØ—z‘z�lØ—z‘z�lô€Gð ˜!•*€AØ€Hr   c                óŽ   € V ^8„  d   QhR\         P                  R\        \        P                  ,          R\         P                  /# )r   ÚlogitsÚmask_npr   )r   r   r
   ÚnpÚint32)r   s   "r   r   r   6   s4   € ÷ ñ ¤§¡ð Ü!(¬¯©Õ!2ðÜ79·x±xñr   c                 ó   € \         P                  ! V4      p\        V P                  4      ^8X  d   \         P                  ! V ^ R7      p \        VP                  4      ^8X  d   \         P                  ! V^ R7      pVP
                  \         P                  8X  g   Q R4       h\        V P                  4      ^8X  g   Q R4       hV P                  w  r4VP                  w  rVW58X  g   Q R4       h\        \         P                  \        W4      4      pV# )r)   )ÚaxiszMask must be int32zLogits must be 2DzBatch size mismatch)	r   r   Úlenr-   Úexpand_dimsr   r6   r   r1   )r3   r4   r   ÚbatchÚvocabÚm_batchÚm_vocabÚrs   &&      r   Úapply_token_bitmaskr@   6   sÎ   € ä�8Š8�GÓ€DÜ
ˆ6�<‰<Ó˜AÔÜ—’ ¨QÔ/ˆÜ
ˆ4�:‰:ƒ˜!ÔÜ�~Š~˜d¨Ô+ˆØ�:‰:œŸ™Ô!Ð7Ð#7Ó7Ð!Üˆv�|‰|Ó Ô!Ð6Ð#6Ó6Ð!Ø—<‘<�L€EØ—z‘zÑ€GØÔÐ2Ð2Ó2ÐÜŒR�X‰XÔ1°&Ó?Ó@€AØ€Hr   )Útypingr   r   r   Únumpyr5   Úmlx.coreÚcorer   Ú_libr   r   r   r	   Únumpy.typingr
   Úcustom_functionr1   r@   © r   r   Ú<module>rI      s?   ð÷ %Ñ $Û Ý Ý ß UÑ UÝ  ð ×Ñô'ó ð'÷Tr   