o
    پiA                  
   @   s  d dl Z d dlZd dlmZmZmZmZmZ d dlZd dlm	Z	 d dl
mZ d dlmZ d dlmZ d dlmZ d dlmZ d d	lmZ d d
lmZ d dlmZ d dlmZmZ d dlmZm Z m!Z! d dl"m#Z# d dl$m%Z%m&Z&m'Z' d dl(m)Z) d dl*m+Z+ d dl,m-Z- d dl.m/Z/ d dl0m1Z1 d dl2m3Z3 d dl4m5Z5m6Z6 d dl7m8Z8 d dl9m:Z: d dl;m<Z<m=Z= d dl>m?Z?m@Z@ d dlAmBZB d dlCmDZDmEZEmFZFmGZGmHZHmIZImJZJmKZK d dlLmMZM eNeOZPeH ZQeI ZReG ZSeF ZTd dlUZUd dlVmWZX eUjYdeXjZdeXjZd eXjZd!eXjZfd"d#Z[d$d% Z\G d&d' d'e	j]Z^G d(d) d)e	j]Z_G d*d+ d+e	j]Z`e`e_d,ZaG d-d. d.e	j]ZbG d/d0 d0e jcZdG d1d2 d2e	j]ZeeeZfeMd3gd4e d5ejgd3ejgd6ehd7dfd8d9ZidS ):    N)AnyIterableOptionalSetTuple)nn)register_split_op)get_forward_context)Qwen3NextConfig)get_pp_group)'get_global_expert_distribution_recorder)ModelConfigForExpertLocation)RMSNorm)mamba_v2_sharded_weight_loader)LayerCommunicatorLayerScatterModes)get_attention_tp_rankget_attention_tp_sizeis_dp_attention_enabled)GemmaRMSNorm)ColumnParallelLinearQKVParallelLinearRowParallelLinear)LogitsProcessor)FusedMoE)QuantizationConfig)RadixAttention)RadixLinearAttention)get_rope)ParallelLMHeadVocabParallelEmbedding)get_is_capture_mode)ForwardBatch)default_weight_loadersharded_weight_loader)Qwen2MoeMLPQwen2MoeSparseMoeBlock)get_global_server_args)	LazyValue
add_prefixcpu_has_amx_supportis_cpuis_cudais_npumake_layersset_weight_attrs)register_custom_opNUM_HEADS_QKNUM_HEADS_VHEAD_QKHEAD_Vc
           "      C   s  t dt d}
}|d || |	 d  }|| d }|d || |	  }|}||
| |  ||  t d| }|| }||
| |  ||  t || }||| |	  }||
| |  ||  t || }||| |	  }||
| |  ||  t || }| |
| |  ||  t d| }| |
| |  ||  ||  t d| }| |
| |  || d  ||	 | |  t d|	| |  }||
| |	  ||	 | |  t d|	| |  }t |t | t |t | t |t | t |t | || }|||  }t |D ](}||
| |  ||  | }||
|  || |  | }t |t | qt ||D ]*}||
| |  ||  | } ||
|  || |  ||  }!t |!t |  qId S )Nr         )tl
program_idarangestoreloadstatic_range)"	mixed_qkvzba
mixed_qkvzmixed_bar1   r2   r3   r4   i_bsi_qk
QKVZ_DIM_TBA_DIM_T	QKV_DIM_Tq_end	blk_q_ptrk_end	blk_k_ptrv_end	blk_v_ptrz_end	blk_z_ptrblk_q_st_ptrblk_k_st_ptrblk_v_st_ptrblk_z_st_ptrb_enda_endi	blk_b_ptrblk_b_st_ptr	blk_a_ptrblk_a_st_ptr r[   P/home/ubuntu/.local/lib/python3.10/site-packages/sglang/srt/models/qwen3_next.py%fused_qkvzba_split_reshape_cat_kernelE   s   














r]   c                 C   s   | j d d}}|| d ||  }tj|| |g| j| jd}	tj|| ||g| j| jd}
tj|| |g|j|jd}t|}|| |f}t| |	|
||| |||||ddd |	|
||fS )Nr   r5   r6   )dtypedevice   )	num_warps
num_stages)shapetorchemptyr^   r_   
empty_liker]   )rA   rB   num_heads_qknum_heads_vhead_qkhead_vbatchseq_len	qkv_dim_tr=   r>   r?   r@   gridr[   r[   r\   fused_qkvzba_split_reshape_cat   sD   


ro   c                       sv   e Zd Z			ddededee deejj	 de
ddf fd	d
Zdd ZdejfddZdejdefddZ  ZS )Qwen3GatedDeltaNetN configlayer_idquant_config
alt_streamprefixreturnc           
         s  t    || _t | _t | _|j| _ts|j	n|j
| _ts"|jn|j| _|j| _|j| _| j| j | _| j| j | _|| _|j| _|| _|j| _|j| _| jd | j | _t| j| jdd | j| jtd|d| _ | j j!j"#d| j j!_"| jd | jd  }| jd }t| j|d|| j| jtd|d| _$t| j|d|| j| jtd|d| _%| jddf}| jddf}	t&| j j!d	 t'| j j!d	t(|||	g| j| ji t)*t+,| j| j | _-t)*t+j,| j| j t+j.d
| _/t'| j/d	t0di t'| j-d	t0di t1| j| jd dt+2 3 |j4d| _5t6| j| jd|dd| j| jtd|d	| _7t8|| j| j | j| j | j| j | j| j| j| j j!9d| j j:| j| j/| j-d| _;d S )Nr6   Fconv1d)
input_sizeoutput_sizebiasrt   tp_ranktp_sizerv   r5   in_proj_qkvz
in_proj_bar   weight_loader)r^   T)eps
group_sizenorm_before_gater_   r^   out_proj)r{   rt   input_is_parallelreduce_resultsr|   r}   rv   )rs   num_q_headsnum_k_headsnum_v_heads
head_q_dim
head_k_dim
head_v_dimconv_weightsr{   
activationA_logdt_bias)<super__init__rr   r   attn_tp_rankr   attn_tp_sizehidden_size_is_cpulinear_num_value_headslinear_num_value_heads_cpur   linear_num_key_headslinear_num_key_heads_cpur   linear_key_head_dimr   linear_value_head_dimr   key_dim	value_dimru   linear_conv_kernel_dimconv_kernel_sizers   
hidden_actr   rms_norm_epslayer_norm_epsilonconv_dimr   r)   rx   weightdata	unsqueezer~   r   delattrr/   r   r   	Parameterrd   zerosr   float32r   r$   RMSNormGatedget_device_modulecurrent_devicetorch_dtypenormr   r   r   squeezer{   attn)
selfrr   rs   rt   ru   rv   projection_size_qkvzprojection_size_baquery_key_settingsvalue_settings	__class__r[   r\   r      s   
	
	

	


zQwen3GatedDeltaNet.__init__c                 C   s`  |  dd | j| j | j| j | j| j | j | j  f }|  dd | j| j d| j | j f }|j| }|j| }| j| j| j| j | j | j| j | j g}| j| j | j| j g}tj||dd\}}}	}
tj||dd\}}|		|	 dd| j}	|
	|
 dd| j}
|	| d| j| j }|	| d| j| j }|||	|
||fS )zQ
        Derives `query`, `key` and `value` tensors from `mixed_qkvzba`.
        Nr6   dimr   )
sizer   r   r   r   r   viewrd   splitreshape)r   rA   rB   new_tensor_shape_qkvznew_tensor_shape_basplit_arg_list_qkvzsplit_arg_list_baquerykeyvaluer>   r?   r@   r[   r[   r\   fix_query_key_value_orderingN  sB   







z/Qwen3GatedDeltaNet.fix_query_key_value_orderinghidden_statesc                 C   s   t stst jrd}nd}|j\}}||k rX| jd urXt rXtj	 }| j
| | |\}}tj| j | |\}}W d    n1 sIw   Y  |
| j ||fS | |\}}| |\}}||fS )Nr   i   )r   _is_npur'   enable_piecewise_cuda_graphrc   ru   r!   rd   cudacurrent_streamwait_streamr~   streamr   )r   r   DUAL_STREAM_TOKEN_THRESHOLDrl   _r   projected_states_qkvzprojected_states_bar[   r[   r\   _forward_input_proj|  s&   


z&Qwen3GatedDeltaNet._forward_input_projforward_batchc                 C   s  |j  }| |\}}| j| j dv r4|r4ts4t||t| j| j	t| j| j	| j
| j\}}}}	nBtrStrStjj||| j| j	 | j| j	 | j
| j\}}}}	n#| ||\}
}}}}}	tdd |
||f\}
}}tj|
||fdd}| j|||	|d}|j}|d|jd }|d|jd }|j|jkrt|}||d |jd d d f< |}| ||}||}|jg |jd d dR  }| |\}}|S )	N)r5   r6      c                 S   s   |  | jd dS )Nr   r   )r   rc   )xr[   r[   r\   <lambda>  s    z,Qwen3GatedDeltaNet.forward.<locals>.<lambda>r   r   )r=   r@   r?   r   )forward_modeis_cuda_graphr   r   r   r   ro   tritoncdivr   r   r   _is_amx_availablerd   ops
sgl_kernel"fused_qkvzba_split_reshape_cat_cpur   mapcatr   rc   r   
zeros_liker   r   )r   r   r   r   r   r   r=   r>   r?   r@   r   r   r   core_attn_out
z_shape_ogcore_attn_out_padoutputr   r[   r[   r\   forward  sh   





zQwen3GatedDeltaNet.forward)NNrq   )__name__
__module____qualname__r
   intr   r   rd   r   Streamstrr   r   Tensorr   r"   r   __classcell__r[   r[   r   r\   rp      s2    
 .rp   c                       sv   e Zd Z			ddededee dedeej	j
 ddf fd	d
Z	ddejdeej deeej  fddZ  ZS )Qwen3HybridLinearDecoderLayerNrq   rr   rs   rt   rv   ru   rw   c                    s   t    || _t|||||| _d| _d}d}|| _tj||j	| j||d| _
| jr=t||||td|ddd| _nt|j|j|j|td|ddd| _t|j|jd| _t|j|jd| _t| j
| j| jdd	| _d S )
NTrs   
num_layersis_layer_sparseis_previous_layer_sparseis_next_layer_sparsemlpz.linear_attnrq   rs   rr   rt   ru   rv   r   intermediate_sizer   rt   rv   r   layer_scatter_modesinput_layernormpost_attention_layernormallow_reduce_scatter)r   r   rr   rp   linear_attnr   rs   r   init_newnum_hidden_layersr   r&   r)   replacer   r%   r   r   r   r   r   r   r   r   layer_communicatorr   rr   rs   rt   rv   ru   r   r   r   r[   r\   r     sR   


z&Qwen3HybridLinearDecoderLayer.__init__r   residualcaptured_last_layer_outputsc                 K   s   | dd }| jj||||d\}}|j s| ||}| j|||\}}| j|}| |||}| j	|||\}}||fS )Nr   r	  )
getr  +prepare_attn_and_capture_last_layer_outputsr   is_idler  prepare_mlpshould_use_reduce_scatterr   postprocess_layer)r   r   r  r	  kwargsr   use_reduce_scatterr[   r[   r\   r     s0   
	z%Qwen3HybridLinearDecoderLayer.forwardNrq   NN)r   r   r   r
   r   r   r   r   rd   r   r   r   r   listr   r   r[   r[   r   r\   r     s2    
;r   c                       s   e Zd Z			ddededee dedeej	j
 ddf fd	d
Zdejdejdeejejf fddZdejdejdedejfddZ	ddejdejdeej dedeeej  defddZ  ZS ) Qwen3HybridAttentionDecoderLayerNrq   rr   rs   rt   rv   ru   rw   c                    s  t    || _|j| _t | _t | _|j| _	| j	| j dks"J | j	| j | _
|j| _| j| jkr>| j| j dks=J n
| j| j dksHJ td| j| j | _|jpZ| j| j
 | _| j
| j | _| j| j | _| jd | _t|dd| _t|dd| _d|v rt|dd | _nt|d	d | _|j| _|| _t|d
d| _| jrtd t| j| j| j| j| j| jdt d| _t |j| j| j	d| j  | jd|| j| jt!d|d	| _"t#| j	| j |jd|d| j| jt!d|d| _$t%| j
| j| j| j||| dd| _&d| _'d}d}t(j)||j*| j'||d| _+| j'r,t,||||t!d|-ddd| _.nt/|j|j0|j1|t!d|-ddd| _.t2|j|j3d| _4t2|j|j3d| _5t2| j|j3d| _6t2| j|j3d| _7t8| j+| j4| j5dd| _9|| _:d S )Nr   r5   g      
rope_thetai'  max_position_embeddingsi    rope_parametersrope_scalingattn_output_gateTzusing attn output gate!)	head_size
rotary_dimmax_positionr  basepartial_rotary_factoris_neox_styler^   Fqkv_proj)r{   rt   r|   r}   rv   o_proj)r{   rt   r   r|   r}   rv   z.attn)num_kv_headsrs   rt   rv   r   r   
.self_attnrq   r   r   r   r   );r   r   rr   r   r   r   r   r   num_attention_headstotal_num_heads	num_headsnum_key_value_headstotal_num_kv_headsmaxr$  head_dimq_sizekv_sizescalinggetattrr  r  r  r   rs   r  loggerwarning_oncer   rd   get_default_dtype
rotary_embr   r)   r"  r   r#  r   r   r   r   r  r  r   r&   r  r   r%   r   r   r   r   r   r   q_normk_normr   r  ru   r  r   r[   r\   r   ;  s   




z)Qwen3HybridAttentionDecoderLayer.__init__qkc                 C   s   | j d urIt rItj }| j | |d| j}| |}tj	| j  |d| j}| 
|}W d    n1 s=w   Y  || j  n|d| j}| |}|d| j}| 
|}||j}||j}||fS )Nr   )ru   r!   rd   r   r   r   r   r,  r5  r   r6  r   rc   )r   r7  r8  r   	q_by_head	k_by_headr[   r[   r\   _apply_qk_norm  s"   



z/Qwen3HybridAttentionDecoderLayer._apply_qk_norm	positionsr   r   c                 C   s  |  |\}}| jrO|j| jd | j| jgdd\}}}|jd d }	|jg |	| jdR  }tj	|ddd\}
}|
j
g |	dR  }
|j
g |	dR  }n|j| j| j| jgdd\}
}}| |
|\}
}| ||
|\}
}| |
|||}| jrt|}|| }| |\}}|S )Nr6   r   r   )r"  r  r   r-  r.  rc   r   r(  rd   chunkr   r;  r4  r   sigmoidr#  )r   r<  r   r   qkvr   q_gater8  v
orig_shaper7  gateattn_outputr   r[   r[   r\   self_attention  s&    
z/Qwen3HybridAttentionDecoderLayer.self_attentionr  r	  r  c                 K   s|   | j j||||d\}}|j s| j|||d}| j |||\}}| j |}| |||}| j |||\}}||fS )Nr
  )r<  r   r   )	r  r  r   r  rE  r  r  r   r  )r   r<  r   r  r   r	  r  r  r[   r[   r\   r     s0   

	z(Qwen3HybridAttentionDecoderLayer.forwardr  r  )r   r   r   r
   r   r   r   r   rd   r   r   r   r   r   r;  r"   rE  r  r   r   r   r[   r[   r   r\   r  9  s`    
 

'r  )	attentionlinear_attentionc                       sx   e Zd Z		ddedee deddf fddZd	ee	 fd
dZ
	ddejdejdedeej dejf
ddZ  ZS )Qwen3NextModelNrq   rr   rt   rv   rw   c                    s   t    | _trtj nd  tjj	jt
 d| _dtdtf fdd}tj|| dd| _tj	jd| _d	| _g | _d S )
N)org_num_embeddingsuse_attn_tp_groupidxrv   c                    sD   t j|   }j|  dkrtd|}ntd|}|| | dS )NrF  	self_attnr  )rt   rv   ru   )ALL_DECODER_LAYER_TYPESlayers_block_typer)   )rK  rv   layer_classru   rr   rt   r[   r\   	get_layer9  s   
z*Qwen3NextModel.__init__.<locals>.get_layerz.layersrv   r   r   )r   r   rr   _is_cudard   r   r   r    
vocab_sizer   r   embed_tokensr   r   r.   r  layersr   r   r   infer_countlayers_to_capture)r   rr   rt   rv   rQ  r   rP  r\   r   '  s    

zQwen3NextModel.__init__rX  c                 C   s(   || _ | j D ]}t| j| dd qd S )N_is_layer_to_captureT)rX  setattrrV  )r   rX  rs   r[   r[   r\   set_eagle3_layers_to_captureQ  s   
z+Qwen3NextModel.set_eagle3_layers_to_capture	input_idsr<  r   inputs_embedsc                 C   s   |d ur|}n|  |}d }g }tt| jD ]1}| j| }	t | |	|||||t|	ddr3|nd d\}}W d    n1 sCw   Y  q|j s`|d u rX| 	|}n| 	||\}}
t|dkrh|S ||fS )NrY  F)rs   r<  r   r  r   r	  r   )
rU  rangelenrV  r   with_current_layerr0  r   r  r   )r   r\  r<  r   r]  r   r  aux_hidden_statesrV   layerr   r[   r[   r\   r   V  s6   



zQwen3NextModel.forwardNrq   r  )r   r   r   r
   r   r   r   r   r  r   r[  rd   r   r"   r   r   r[   r[   r   r\   rH  &  s2    *rH  c                   @   s   e Zd ZdZdZdZdZdS )HybridLayerTyperF  swa_attentionrG  mambaN)r   r   r   full_attentionre  rG  mamba2r[   r[   r[   r\   rd    s
    rd  c                       s   e Zd ZdZg dddgdZ		d'ded	ee d
eddf fddZ	e
dd Ze 	d(dejdejdedeej fddZdd Zdd Zdd Zdd Z	d)deeeejf  dedee fd d!Zed"d# Zd(d$eee  fd%d&Z  ZS )*Qwen3NextForCausalLMF)q_projk_projv_proj	gate_projup_proj)r"  gate_up_projNrq   rr   rt   rv   rw   c                    s   t    | _t  _ jjr jjsJ |d ur#t|dr# j|_| _	t
||td|d _t|j|j||jtd|t jd _t| _d _t fdd _d S )	Npacked_modules_mappingmodelrR  lm_head)rt   rI  rv   rJ  Fc                      s   dd t  jjD S )Nc                 S   s(   i | ]\}}t |jtr||j qS r[   )
isinstancer   r&   get_moe_weights).0rs   rb  r[   r[   r\   
<dictcomp>  s    

zCQwen3NextForCausalLM.__init__.<locals>.<lambda>.<locals>.<dictcomp>)	enumeraterq  rV  r[   r   r[   r\   r     s    
z/Qwen3NextForCausalLM.__init__.<locals>.<lambda>)r   r   rr   r   pp_groupis_first_rankis_last_rankhasattrrp  rt   rH  r)   rq  r   rT  r   r'   enable_dp_lm_headrr  r   logits_processorcapture_aux_hidden_statesr(    _routed_experts_weights_of_layer)r   rr   rt   rv   r   rx  r\   r     s.   



zQwen3NextForCausalLM.__init__c                 C   s   | j jS r  )r  r   rx  r[   r[   r\   routed_experts_weights_of_layer  s   z4Qwen3NextForCausalLM.routed_experts_weights_of_layerr\  r<  r   r]  c                 K   s6   |  ||||}d }| jr|\}}| ||| j||S r  )rq  r  r~  rr  )r   r\  r<  r   r]  r  r   ra  r[   r[   r\   r     s   	zQwen3NextForCausalLM.forwardc                 C   s   | j jj| jjfS r  )rq  rU  r   rr  rx  r[   r[   r\   get_embed_and_head  s   z'Qwen3NextForCausalLM.get_embed_and_headc                 C   s8   | j j`| j`|| j j_|| j_tj  tj  d S r  )rq  rU  r   rr  rd   r   empty_cachesynchronize)r   embedheadr[   r[   r\   set_embed_and_head  s   

z'Qwen3NextForCausalLM.set_embed_and_headc                 C   s
   | j jjS r  )rq  rU  r   rx  r[   r[   r\   	get_embed  s   
zQwen3NextForCausalLM.get_embedc                 C   sJ   t | jdr| jj| jjkrd S | jj`|| jj_tj	  tj
  d S )Ntarget_hidden_size)r|  rr   r  r   rq  rU  r   rd   r   r  r  )r   r  r[   r[   r\   	set_embed  s   


zQwen3NextForCausalLM.set_embedweightsis_mtpc              	   C   s  g d}t jddd| jjd}t|  }t }|D ]\}}|r6d|vr%q|dv r0|dd	}n|dd
}|s=d|v r=qd|v rBqd|v rL|dd	}|drX|dd}n|drc|dd}|D ]6\}	}
}|
|vroqed|v rtqe||
|	}|dr||vrqe||vrqe|}|| }t	|d}||||  nS|D ]7}|\}	}
}}|
|vrq||
|	}|ds|dr||vrq|}|| }t	|d}||||||d  n|dr||vrq|| }t	|dt
}||| || q|S )N))r"  rj  r7  )r"  rk  r8  )r"  rl  rA  )ro  rm  r   )ro  rn  r5   rm  	down_projrn  )ckpt_gate_proj_nameckpt_down_proj_nameckpt_up_proj_namenum_expertsmtp)zmtp.fc.weightz mtp.pre_fc_norm_embedding.weightzmtp.pre_fc_norm_hidden.weightzmtp.rq   rq  zrotary_emb.inv_freqz.self_attn.r%  z.k_proj.k_scalez.attn.k_scalez.v_proj.v_scalez.attn.v_scalezmlp.expertsz.biasr   _bias)shard_id	expert_id)r   make_expert_params_mappingrr   r  dictnamed_parameterssetr  endswithr0  r#   add)r   r  r  stacked_params_mappingexpert_params_mappingparams_dictloaded_paramsnameloaded_weight
param_nameweight_namer  replaced_nameparamr   mappingr  r[   r[   r\   load_weights  s   




z!Qwen3NextForCausalLM.load_weightsc                 C   s   t |j|jd dS )N)r   num_logical_experts
num_groups)r   r  r  )clsrr   r[   r[   r\   $get_model_config_for_expert_locationk  s
   z9Qwen3NextForCausalLM.get_model_config_for_expert_location	layer_idsc                 C   sZ   | j jsd S d| _|d u r | jj}| jd|d |d g d S | jdd |D  d S )NTr6   r`   c                 S   s   g | ]}|d  qS )r5   r[   )ru  valr[   r[   r\   
<listcomp>  s    zEQwen3NextForCausalLM.set_eagle3_layers_to_capture.<locals>.<listcomp>)ry  r{  r  rr   r  rq  r[  )r   r  r   r[   r[   r\   r[  s  s   z1Qwen3NextForCausalLM.set_eagle3_layers_to_capturerc  r  )F) r   r   r   fall_back_to_pt_during_loadrp  r
   r   r   r   r   propertyr  rd   no_gradr   r"   r   r  r  r  r  r   r   boolr   r  classmethodr  r  r   r[  r   r[   r[   r   r\   ri    sX    *

x
 ri  r   )mutates_argsr   rs   rw   c                 C   sh   t  }|j}|j}|| }|| |}| | ks)J d|  d|  ||j| d S )Nz Output tensor element mismatch: z != )r	   r   attention_layers_forwardnumelr   rc   copy_)r   r   rs   contextr   r  attention_layerretr[   r[   r\   gdn_with_output  s   r  )jenumloggingtypingr   r   r   r   r   rd   r   )sglang.srt.compilation.compilation_configr   0sglang.srt.compilation.piecewise_context_managerr	   sglang.srt.configs.qwen3_nextr
   sglang.srt.distributedr   #sglang.srt.eplb.expert_distributionr   sglang.srt.eplb.expert_locationr   /sglang.srt.layers.attention.fla.layernorm_gatedr   r   'sglang.srt.layers.attention.mamba.mambar   sglang.srt.layers.communicatorr   r   sglang.srt.layers.dp_attentionr   r   r   sglang.srt.layers.layernormr   sglang.srt.layers.linearr   r   r   "sglang.srt.layers.logits_processorr   ,sglang.srt.layers.moe.fused_moe_triton.layerr   *sglang.srt.layers.quantization.base_configr   !sglang.srt.layers.radix_attentionr   (sglang.srt.layers.radix_linear_attentionr   "sglang.srt.layers.rotary_embeddingr   *sglang.srt.layers.vocab_parallel_embeddingr   r    +sglang.srt.model_executor.cuda_graph_runnerr!   ,sglang.srt.model_executor.forward_batch_infor"   $sglang.srt.model_loader.weight_utilsr#   r$   sglang.srt.models.qwen2_moer%   r&   sglang.srt.server_argsr'   sglang.srt.utilsr(   r)   r*   r+   r,   r-   r.   r/   sglang.srt.utils.custom_opr0   	getLoggerr   r1  rS  r   r   r   r   triton.languagelanguager7   jit	constexprr]   ro   Modulerp   r   r  rM  rH  Enumrd  ri  
EntryClassr   r   r  r[   r[   r[   r\   <module>   s    (

	
X,  a i_ z
