o
    پi                     @   s:   d dl Z d dlmZ e jdejdejdejfddZdS )    NAPPLY_SOFTMAXN_EXPTS_ACTN_EXPTS_PADc                 C   s  t d}|	d urt |	}||krd S | || 7 } ||| 7 }||| 7 }||| 7 }t d|}t d|}||
k }t | | }t || }|t j}t |}t || }|t j}t || d}t j|| d|d t 	  |r|||  }n|}t || | d S )Nr   )mask)
tl
program_idloadarangetofloat32softmaxsumstoredebug_barrier)Yi	stride_ymDY
stride_dymX	stride_xmDX
stride_dxmn_rowsNRowsn_expts_totr   r   r   pid_moffs_xnoffs_ynmask_xny_indxxydysdx r%   ^/home/ubuntu/.local/lib/python3.10/site-packages/triton_kernels/topk_details/_topk_backward.py_topk_backward   s2   


r'   )tritontriton.languagelanguager   jit	constexprr'   r%   r%   r%   r&   <module>   s    