o Ti:,@sddlZddlZddlmZddZddZedddiedd diejd d Z edd diedd diejddZ Gdddej j Z GdddZdS)NcCsP|d8}||d?O}||d?O}||d?O}||d?O}||d?O}|d7}|S)NnrrZ/home/ubuntu/.local/lib/python3.10/site-packages/deepspeed/ops/sparse_attention/softmax.pynext_power_of_2s     r cCs|dkrdS|dkr dSdS)Nirirrrrrrr num_warpss r cOt|d|dSNBLOCKr argsmetarrr "rTNcOr rr rrrr r#rc 'Ks| d}| d}td}td}||}||}td||}td||}||d}t|d}t|d}||k}t|||d}t|||dd}t|||dd}t|||dd}t|||dd}|||||||||}tj||td d } | tj} | d r| |} | d r||||| ||||| || |}!tj|!|dd }"| |"} | d r||| |||}#tj|#|td d }$| d rt|$dktd d}$| |$} | dr<|||||| || |}%tj|%|td d }&| dr8t|&dktd d}&| |&} t| } tj || |ddS)Nrrrrrrinfmaskother APPLY_SCALE APPLY_RPE APPLY_KP_MASK KP_MASK_MULgAPPLY_ATTN_MASK ATTN_MASK_MULr) tl program_idarangeloadwherefloattofloat32softmaxstore)'XscaleLUTRPEKP_MATTN_Msizemax stride_zx stride_zrpe stride_hrpe stride_srpe stride_zkpm stride_zattnmrrrpidhmpidzrxmrbmrxnrbnheadersizeoffsetcheckrbmnblockidcolumnidrowidheadidpxxprperpepkp_mkp_mpattn_mattn_mrrr _forward"sP   $ 4   $  rScOr Nrrrrrrr rZrcOst|d|dSrTrrrrr r[rcKs\td}td} |d} |d} || } || } td| | }td| | }|| d}t|d}t|d}||k}t|||d}t|||d}|| ||| | | | |}|| ||| | | | |}tj||dd}tj||dd}|tj}|tj}||t||d|}tj|||ddS) Nrrrrrrrr$) r%r&r'r(r)r+r,sumr.)r/r0DXr1r5r6 stride_zdxrr<r=rrr>r?r@rArBrCrDrErFrGrLdxyrrr _backwardZs,   $$  rZc@s6eZdZeZeddZeddZeddZdS)_sparse_softmaxc Cs<tjgtj|jd}|}t|jdD]}t|||ddddfdf}qt |}tj |dddd|dd<t |}| dddf}| dddf} | dddf} tj || | |fddd} |dd|}tj ||fddd} t| | ftj|} | t|fS)Ndtypedevicer)dimrrr)torchtensorint64r^clonerangeshapecatrU zeros_likecumsumr'nonzerostackviewnumeltypeint32r+intmax)layoutblockr^_emptysizeshoffsetsidxheadrowscolumnscorerBlutrrr make_luts( z_sparse_softmax.make_lutcs|dkrdnd}|durd}d\}}}tjd|j|jd}nd}|d|d|d}}}|durDd}d}tjd|j|jd}nd}|d}|dur^d}d}tjd|j|jd}nd}|d}|jd|||||d k|d kd }fd d }t|||| |||| |d|||||f i|||||| |_ |_ | |_ ||_ ||_ ||_||_||_||_||_|S) N?FT)rrrrr\rrmul)rrrr r"r!r#csddgSNrrroptMrsspdimsrr rsz)_sparse_softmax.forward..)raemptyr]r^striderfrS mark_dirtysave_for_backwardrrsmaxlutr0 apply_scale apply_rpe apply_kp_maskapply_attn_mask kp_mask_modeattn_mask_mode)ctxrLr0rNkey_padding_mask attn_maskrrrrsr} num_blocksrbenchtimerrr7r8r9rr:rr;rgridrrr forwards^ "       z_sparse_softmax.forwardcsrj\}}|jdfdd}t||j||j|d|djd|ddddddddddddddfS)NrcsjdjdjgSr)rrsrrrrr rsz*_sparse_softmax.backward..)r) saved_tensorsrfrZr0rrrs)rrXrLr}rrrr backwards  ."z_sparse_softmax.backwardN) __name__ __module__ __qualname__dict bwd_kernels staticmethodr~rrrrrr r[{s  >r[c@s@eZdZdZddZddZdddZ dd d Zd S)SoftmaxarBlock-Sparse Softmax class; this class computes softmax on a block sparse matrix. It is also able to apply either/all of the following masks: - relative position embedding - key padding mask - attention mask For more details about sparsity config, please see `Generative Modeling with Sparse Transformers`: https://arxiv.org/abs/1904.10509 cOstj|i|S)N)r[apply)rkwargsrrr sparse_softmaxszSoftmax.sparse_softmaxcCs2|f}||jvrt|j|j||j|<|j|S)zCGenerates the sparsity layout used in block-sparse softmax ) lut_cacher[r~rrrs)selfr^keyrrr r~s  zSoftmax.make_lutFcCs4||_|j|_||_||_||_t|_ dS)aInitialize the Block-Sparse Softmax class. Arguments: layout: required: sparsity layout tensor block: required: an integer determining the block size. bench: optional: set if you want to do benchmarking N) rUitemrrfrrrrsrrr)rrrrsrrrr __init__s   zSoftmax.__init__rNaddc Csdg}|dur|j|jkrtd|j|dur%|j|jkr%td|j|dur6|j|jkr6td|j||j\} } t||||||||j|j| |j| |j |}|d|_ |S)aApplies softmax on a Block-Sparse input tensor. For more details about sparsity config, please see `Generative Modeling with Sparse Transformers`: https://arxiv.org/abs/1904.10509 Arguments: x: required: a block-sparse tensor that softmax is applied on it; computation will be in place and result will be returned in the same tensor scale: optional: a float value; x values will be multiplied by this value before normalization. Default value is 1.0. rpe: optional: a tensor same dimension as x that is used as relative position embedding key_padding_mask: optional: a mask tensor of size (BatchSize X SequenceLength) attn_mask: optional: a mask tensor of size (SequenceLength X SequenceLength); currently only 2D is supported key_padding_mask_mode: optional: a boolean determining if key_padding_mask needs to be added or multiplied attn_mask_mode: optional: a boolean determining if attn_mask needs to be added or multiplied Return: x: a block-sparse tensor contains normalized input x using softmax; and masks applied if given Nz&relative position embedding must be %szAttention mask must be %szKey padding mask must be %sr) r] ValueErrorr~r^rrrrsrrtime_y) rrLr0rNrrkey_padding_mask_moderrr}rrrr __call__s zSoftmax.__call__)F)rNNNrr)rrr__doc__rr~rrrrrr rs r)ratritontriton.languagelanguager%r r heuristicsjitrSrZautogradFunctionr[rrrrr s  5e