*`i ddlmZmZddlZ ddlZddlmZn#e$rZ ede dZ [ wwxYwej dej dej fdZ d dej dej d eed eeefd ZdS) )ListOptionalNzTriton is not installedNUM_SMS BLOCK_SIZEc tjd} tj||} tj| || z|D]} | | z} | | z|z} || ntj|| z}| tjd|z}| dztjd|dzz}||k}||k}tj|||zz|z|}|dddftjdddddfz dzdk}||}tj|||zz|ztd ||zdS)aApply a bitmask to logits in-place using Triton. The bitmask is a 01 bitwise compressed tensor, where 0 means the token is masked and 1 means the token is not masked. After applying the bitmask, the masked logits will be set to -inf. Parameters ---------- logits_ptr : tl.tensor Pointer to the logits tensor to apply the bitmask to. bitmask_ptr : tl.tensor Pointer to the bitmask tensor to apply. indices_ptr : Optional[tl.tensor] Optional pointer to indices tensor specifying which rows to apply the mask to. num_rows : int Number of rows to process. If indices_ptr is provided, this is the number of unique indices. vocab_size : int Size of the vocabulary dimension. If the logits does not have a vocab padding, this is the same as the logits's second dimension. Otherwise, this is the actual size of the vocabulary. logits_strides : int Stride between rows in the logits tensor. bitmask_strides : int Stride between rows in the bitmask tensor. NUM_SMS : int Number of streaming multiprocessors to use. BLOCK_SIZE : int Size of processing blocks. rN inf) tl program_idcdivrangeloadarangereshapestorefloat) logits_ptr bitmask_ptr indices_ptrnum_rows vocab_sizelogits_stridesbitmask_stridesrrpid num_blockswork_idrow_id block_offsetbatch_idoffsetsbitmask_offsets vocab_maskpacked_bitmask_maskpacked_bitmaskbitmasks /home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/xgrammar/kernels/apply_token_bitmask_inplace_triton.py"apply_token_bitmask_inplace_kernelr( s^ -  CZ00J8CJ!6@@  J&*, : (066bgkF>R6S6S1j!9!99&",ryJ">> .. D.II :D&t,  )bF4G4G4I4I.IJ      r))NN)typingrrr9tritontriton.languagelanguager ImportErrorerrjit constexprr(TensorintrMr)r'rYs!!!!!!!! :MMM       ::: +/ 0 0c9:@ \@  @ @ @  @ L!%#' .. L. \. .d3i ......s / */