a zþdý<ã@sÎddlZddlZddlmZe ddd„i¡ejejejejejejdœdd„ƒƒZddd „Z d d „Z ejejejejejejd œd d„ƒZ ddd„Z dd„Z ddd„ZGdd„dejjƒZGdd„dƒZdS)éNÚEVEN_KcCs|d|ddkS)NÚKÚTILE_Kr©)Únargsrrúf/var/www/html/stable-diffusion-webui/venv/lib/python3.9/site-packages/triton/ops/blocksparse/matmul.pyÚór)ÚTILE_MÚTILE_NrÚBLOCKrc*Cs:t d¡|}||d7}t d¡}t |d¡}t |d¡}||t d|¡|}t d|¡}||||||dd…df||ddd…f|}t |d¡}||t d|¡|}t d|¡} ||||||ddd…f| | dd…df| }!tj||ftjd}"t|d| ƒD]Š}#|rPt |¡}$t |!¡}%n@tj||ddd…f|#kdd}$tj|!| dd…df|#kdd}%|"t |$|%¡7}"|||7}|!|| 7}!q0|" |j j ¡}&t d|¡|}'t d|¡|}(||| || |'dd…df| |(ddd…f|})tj |)|&dd dS) Nééér©Údtypeg)ÚmaskÚotherT©r) ÚtlÚ program_idÚloadÚarangeÚzerosÚfloat32ÚrangeÚdotÚtorÚ element_tyÚstore)*ÚAÚBÚCZ stride_zaÚ stride_haZ stride_maÚ stride_akÚ stride_zbÚ stride_hbÚ stride_bkZ stride_nbÚ stride_zcÚ stride_hcZ stride_mcZ stride_ncrZ grid_offsetÚlutr r rr rÚblock_idZoff_zÚoff_hZstart_amÚoffs_amÚoffs_akZa_ptrsZstart_bnÚoffs_bnÚoffs_bkZb_ptrsÚaccÚkÚaÚbÚcÚoffs_cmÚoffs_cnÚpcrrrÚ _sdd_kernelsf   ÿþýü ÿþýü     ÿþýür9c CsÐ| d¡dkr$| d¡dkr$| ¡}| d¡dkrH| d¡dkrH| ¡}|rd||}}| | }}|rldnd} |rxdnd} |j| |j| } } | | kr°td| ›d| ›dƒ‚| durâtj|jd |jd ||f|j|jd }n(| j|jd |jd ||fksJ‚| }d|jd|jd g}t||||| d ¡| d¡| |rJdnd¡| |r\dnd¡| d ¡| d¡| |r~dnd¡| |rdnd¡| d ¡| d¡| d¡| d¡| d |||d |d d d |S)Nrr réþÿÿÿéÿÿÿÿzInner dimension mismatch (A: z vs B: ú)r©rÚdeviceé é)r r rr Ú num_stagesÚ num_warps) ÚstrideÚ contiguousÚshapeÚ ValueErrorÚtorchÚemptyrr>r9)r3r4Útrans_aÚtrans_bÚtrans_cÚspdimsÚblockr*ZwidthsÚoutZa_dimZb_dimZKaZKbr5ÚgridrrrÚ sdd_matmulTs6   *$44  ù rPcCs&|jdd |¡ ¡}| ¡}|dfS)NF©Úas_tuple)ÚnonzerorÚintrD)ÚlayoutrMr>r*rrrÚsdd_lutvsrV)r r rÚ GROUP_SIZE_Mr c4Cst d¡}t d¡}t d¡}t d¡}t |||||¡\}}t d¡}||d}t |d¡}t |d¡}t |d¡}t |d¡} ||}!t |!d¡}"t |"d¡}"t d|¡}#t d|¡}$||||"||#dd…df||$ddd…f|}%||t d|¡}&t t |&||¡|¡}&t |!¡}'t |'d¡}'|'t d|¡}(|||| ||&ddd…f| |(dd…df| })tj||ftj d}*|!d7}!t |!d¡}+t |+d¡}+t |!¡},t |,d¡},t |d| ƒD]Ž}-tj|%dd }.tj|)|&ddd…f|kd }/|*t  |.|/¡7}*|%|+7}%|)|,| 7})|!d7}!t |!d¡}+t |+d¡}+t |!¡},t |,d¡},qê|*  |j j¡}0||t d|¡}1||t d|¡}2|| | || |1dd…df| |2ddd…f|}3tj|3|0|2ddd…f|kd dS) Nrr rr@rérTr)rrZ num_programsZ swizzle2drZ multiple_ofrZmax_contiguousrrrrrrrr)4r r!r"Z stride_azr#Z stride_amr$r%r&r'Z stride_bnr(r)Z stride_cmZ stride_cnZDS0ZDS1r*r r rrWr Zpid_mZpid_nZ num_pid_mZ num_pid_nZpidzÚheaderÚoffsetrÚcolumnr,Zpincr+r-r.Úpar/Zstart_bkr0Zpbr1Zinc_aZinc_br2r3r4r5r6r7r8rrrÚ _dsd_kernel‚s‚          ÿþý   ÿþý      ÿþýür]c sÆ| d¡dkr$| d¡dkr$| ¡}| d¡dkrH| d¡dkrH| ¡}|||rTdnd} | d¡‰| d¡} | |r|dnd¡‰|j} ˆ} | }|r˜ˆn| }|r¤| nˆ}| durÌtj| |||f| |jd}n| j| |||fksâJ‚| }d}‡‡‡fdd„}t||||| d¡| d¡| |r$dnd¡| |r6dnd¡| d¡| d¡| |rXdnd¡| |rjdnd¡| d¡| d¡| |rŒdnd¡| |rždnd¡ˆ| |||t |d ƒ|d d d d |S) Nrr rrr=é€cst ˆ|d¡ˆˆgS)Nr )ÚtritonZcdiv)Úmeta©ÚBS0ZBS3Úwidthrrrær zdsd_matmul..r?r@)r r rr rArBrW) rCrDÚsizerrGrHr>rEr]Úmin)r3r4rIrJrKrLrMr*rcrNZAS1ÚBS1rZCS0ZCS1ZCS2ZCS3r5r rOrrarÚ dsd_matmulÎs:    444ù rgcCs¢t ||rdnd¡}t |¡jdd\}}| ¡}||}|rL|jdd} n| dd¡jdd} |  d¡} t |¡} tj|dd…dd | dd…<t  | | dt | ¡¡} | dd…df|} |   ¡} | dd…| dd…8<||}|   dd¡  d|¡} || dd…dd…f<| dd…df|d|8<| | |dk| | |dkdf<|   d¡} |rvtj | |jd }nštjgtj|jd }d}t| d¡ƒD]r}||dd…dd…f  ¡ ¡}| ¡}dtj ||jd ||dk<t |||j|jdkdf¡}||7}qœ|||}|dd…|dd…||8<|  dd¡  d|¡}|r||dd…dd…f<|dd…df|d|8<n<|||dd…dd…f<|dd…df|d||8<|| |dk|| |dkdf<|  d¡}| d¡}| d|d |} ||}tj| |||fdd   d¡ ¡}tj| |fdd   d¡ ¡}tjd |j|jd}t ||f¡}t ||f¡}| tj¡ |¡}||fS)a Generates the look-up table for incrementing pointers in the DSD/DDS matmul. Example (BLOCK=32, STEP=16) [[1, 0, 0, 1, 0], [0, 1, 1, 0, 1], [1, 0, 1, 0, 0]] Then the offsets for A are [0 , 16, 32, 48] <- row 0 \----/ \----/ col=0 col=3 [64, 80, 96, 112, 128, 144] <- row 1 \----/ \----/ \------/ col=1 col=2 col=3 [160, 176, 192, 208] which leads to increments table [0, 16, 16, 16, || 64, 16, 16, 16, 16, 16, || 160, 16, 16, 16] Because B is dense, the offsets are [0, 16, 96, 112] <- row 0 [32, 48, 64, 80] <- row 1 [0, 16, 64, 80] <- row 2 rr TrQFrNr;)Zdim)r>r=r@é)r>r)rGÚsumZ ones_likerSÚflattenZ transposerdZ zeros_likeZcumsumreÚcloneÚviewÚrepeatrr>ZtensorÚint64rÚlongÚcatÚTÚstackrDrrÚtypeÚint32r)rUrMÚstepZtransr>ÚsizesZhead_idZcol_idÚsegmentsZnnzZ num_blocksÚoffsetsZB_idxZB_incsÚdivZA_idxÚcurrent_offsetÚzZlayoutwZmsumZA_incsrcrYZincsÚpadr*rrrÚdsd_lutôsd     "  $"$    r}c Cs"t||| | | ||||| d S©N)rN)rg) r3r4rIrJrKrLrMr*rcrNrrrÚ dds_matmulZsrc@s0eZdZeeedœZedd„ƒZedd„ƒZ dS)Ú_matmul©ÚsddÚdsdÚddsc Csxtj||||||||| | |d }| ||¡| |_| |_| |_||_||_||_||_ ||_ ||_ ||_ |du|_ |Sr~)r€ÚfnZsave_for_backwardÚda_lutÚda_widthÚdb_lutÚdb_widthÚmoderLrMrIrJrKÚhas_out)Úctxr3r4rIrJrKrŠrLrMÚc_lutÚc_widthr†r‡rˆr‰rNr5rrrÚforwardfs"  z_matmul.forwardc Csè|j\}}d\}}|j}|jdrh|d|d|d}tj||||j|j |j|j|j |j |j ƒ }|jdr¸|d|d|d}tj||||j |j|j|j|j |j |j ƒ }|jrÂ|nd} ||dddddddddddd| fS)N)NNrr r)Z saved_tensorsrŠZneeds_input_gradr€r…rKrJrIrLrMr†r‡rˆr‰r‹) rŒÚdcr3r4ÚdaÚdbrŠZmode_daZmode_dbZdoutrrrÚbackward{s$  "ÿ "ÿ  þz_matmul.backwardN) Ú__name__Ú __module__Ú __qualname__rPrgrr…Ú staticmethodrr“rrrrr€bs   r€c@s eZdZddd„Zddd„ZdS) ÚmatmulFc CsB|dvrtdƒ‚||_||_||_||_||_||_|j|_t |dƒ}|jdkr”t |||ƒ\|_ |_ t |||d|ƒ\|_|_t |||d|ƒ\|_|_|jdkrèt ||||j |ƒ\|_ |_ t |||ƒ\|_|_t ||||j|ƒ\|_|_|jdkr>t ||||j|ƒ\|_ |_ t ||||j |ƒ\|_|_t |||ƒ\|_|_dS) Nrz"Supported modes are: sdd, dsd, ddsr?r‚TFrƒr„)ÚNotImplementedErrorrMrŠrIrJrKrUrErLrerVrrŽr}r†r‡rˆr‰) ÚselfrUrMrŠr>rIrJrKrurrrÚ__init__•s,    zmatmul.__init__NcCsBt |||j|j|j|j|j|j|j|j |j |j |j |j |¡}|S)N)r€ÚapplyrIrJrKrŠrLrMrrŽr†r‡rˆr‰)ršr3r4rNr5rrrÚ__call__­sûzmatmul.__call__)FFF)N)r”r•r–r›rrrrrr˜“s r˜)N)N)N)rGr_Ztriton.languageÚlanguagerZ heuristicsZjitZ constexprr9rPrVr]rgr}rZautogradZFunctionr€r˜rrrrÚs( ÿ ù@ "  ùK &f 1