o پik @sddlZddlmZddlmZmZejdejfddZ ejdejfdd Z ejd d Z ejdejd ejfd dZ ejdejd ejdejdejfddZ ejdejfddZejdejdejdejdejdejf ddZdS)N)_expt_data_compute_expt_data_memsetBLOCK_Nc Cs||d|}t|g|jj}t|D]9}||td|}||k}tj|||d} t| d| |} |t| d7}tj ||| |d||7}qdS)Nrrmask) tlzerosdtype element_tyrangearangeloadcumsumsumstore) ExpertHistFinalExpertOffs hist_sizerloop_iterationsxioffs_nmask_nhist2 tok_startsrc/home/ubuntu/.local/lib/python3.10/site-packages/triton_kernels/routing_details/_routing_compute.py_routing_compute_expt_offss  rBLOCK_Mc Cstd|}d}td||D]7}||||} tj|| ||kd} t| d|} |t| d7}tj|| | | ||kd||7}qdS)Nrr)rr r rrrr) PartialHistshape_pm stride_pm stride_pnrexpt_idoffs_mcurr_sum_offscurroutrrr_routing_compute_indx_offss  r+cCs2d}||@}||@}t||k||||}|S)N)rwhere)rykey_maskkxkyzrrr _keyed_add$s r3 N_EXPTS_ACTc Csrt| tjr| jrt| } | | } t| | dktd| | }|| | |}tj|||| kddtj }|d>|Btj }t |d}|d?}|| | |d@}|dk}tj|||d}|d@d B}t |dt }|d d@}tj||||||d}|tj| ||d7}||7}tj ||||dtj ||||dtj ||||ddS) Nir)rotherirr,r) isinstancertensorr is_ptrr static_assertr touint32sortassociative_scanr3r)pid_m GatherIndx ScatterIndxGateScalExptScalExptIndx PartialOffsr"r# TokensStartn_tokensrr4n_gates local_offsr(expertkv_pairsr gate_scalrexpts_and_inclusive_run_lengthsexclusive_run_lengthsgatesrrr_routing_compute_indx0s,     rQSIZESBLOCKcCsXtd}||krt| | ||||||dS||8}t||||||||||| | | dS)Nr)r program_idrrQ)rArBrCrDrErFr"r#rGrHrr4Hist MDTileStartstile_starts_stridem MDTileInfotile_info_stridemfirst_tile_dim_log2rRrSblocks2apidrrr_combined_routing_computeTs  r]c Cstd}|d}|d}d|>d} td||D]A} | td|} tj|||| || |kd} t| |k| | @| } t| |kd| } tj|||| || | |kdqdS)Nr rr)rrTr r rr-r) Bitmatrix stride_bm stride_bnshape_bncutoffrr@ cutoff_word cutoff_bit cutoff_maskstart_nrvaluesrrr_routing_clear_bitmatrixds  "&riBLOCK_Ac Cstd}||krt||| | ||||dS|||kr%t||||dS|||kr8t|| | | |||dS|||d|td|}||k}tj||||ddS)a3 This kernel essentially combines 6 different pieces of functionality, statically branching on the value of tl.program_id(0) to decide which codepath to take. pid == 0: create the token cumsum 1 <= pid <= SIZES: create a tile cumsum SIZES < pid < blocks1a: initialise MDTileInfo to 0xffffffff blocks1a <= pid < blocks1a + n_expts_tot: compute_indx_offs pid == blocks1a + n_expts_tot: compute_expt_offs pid > blocks1a + n_expts_tot: initialise Indx to sentinel As each of these is a relatively trivial workload, launching them from this single trampoline is beneficial as they can execute on different streaming multiprocesses in parallel. rrrN)rrTrrr+r r)IndxsizesentinelrSrrr n_expts_totr r!r"r#MDStartsrWblocks1arXrZrRrjrrr\r(rrrr_combined_routing_memsetrs    rq)tritontriton.languagelanguager _expt_datarrjit constexprrr+r3rQr]rirqrrrrsN     #