a
    zd                     @   sl  d dl Z d dlZd dlmZ ddlmZmZ dd Zdd Z	ej
ejdd	d
dddddejd	dd
dddddejd	dd
dddddejdd	d
dddddejddd
dddddejddd
dddddejddd
dddddejdd
d
dddddejdd
d
dddddejdd	ddddddejd	dddddddejd	dddddddejdd	ddddddejddddddddejddddddddejddddddddejdd
ddddddejdd
ddddddge	  g deedddeddd iejejejejejejejejdddZG dd de jjZejZdS )    N   )early_config_pruneestimate_matmul_timec                    s    fddS )Nc                    s   |     S N)Zzero_)nargsname Z/var/www/html/stable-diffusion-webui/venv/lib/python3.9/site-packages/triton/ops/matmul.py<lambda>	       zinit_to_zero.<locals>.<lambda>r	   r   r	   r   r
   init_to_zero   s    r   c                  C   s   g } dD ]}dD ]x}dD ]n}dD ]d}|dkr0dnd}|  tj|||dd	||d
 dD ]*}|  tj||||d	||tdd qXq qqq| S )N)               )       )r   @   )r   r         r   r   r   r   BLOCK_MBLOCK_NBLOCK_KSPLIT_K
num_stages	num_warps)r   r      r   C)r   r   Zpre_hook)appendtritonConfigr   )configsr   Zblock_mZblock_kZblock_nr   Zsplit_kr	   r	   r
   get_configs_io_bound   s     
r&   r   r   r   r   r   r    r   r   r   r   r   )MNK
   )r   Z
perf_modelZtop_k)r%   keyZprune_configs_byEVEN_Kc                 C   s   | d | d | d   dkS )Nr)   r   r   r   r	   )argsr	   r	   r
   r   <   r   r   )r   r   r   GROUP_Mr   r,   ACC_TYPEc           &      C   s  t d}t d}|| d | }|| d | }|| }|| }t|||  |}|| ||  }|| | }|| t d| }|| t d| }t t || ||}t t || ||}|| t d| } | |d d d f | | d d d f |   } || d d d f | |d d d f |	   }t j||f|d}!t|d| | D ]}"|rzt | }#t |}$n@t j| | d d d f |"k dd}#t j|| d d d f |"k dd}$|!t 	|#|$7 }!| || | 7 } ||| | 7 }qZ|!
|jj}!|| t d| }|| t d| }||d d d f |
 |d d d f |   }||k d d d f ||k d d d f @ }%|dkrt j||!|%d nt j||!|%d d S )Nr   r   )dtypeg        )maskother)r1   )tlZ
program_idminZarangeZmax_contiguousZmultiple_ofzerosrangeloaddottor0   Z
element_tystoreZ
atomic_add)&ABr!   r'   r(   r)   Z	stride_amZ	stride_akZ	stride_bkZ	stride_bnZ	stride_cmZ	stride_cnr   r   r   r.   r   r,   r/   pidZpid_zZgrid_mZgrid_nwidthZgroup_idZ
group_sizeZpid_mZpid_nZrmZrnZramZrbnZrkacckabr1   r	   r	   r
   _kernel   sD    +

,,
  ,(
rC   c                   @   s,   e Zd ZeZi Zedd Zedd ZdS )_matmulc                    s  | j }| ddkr*| ddkr*|  } |ddkrN|ddkrN| }| jd |jd ksjJ d| j\ }|j\}tj f|| jd}| jtjtjtj	fv rt
j	nt
j} fdd}t| | || || d| d|d|d|d|dd|d |S )	Nr   r   zincompatible dimensions)devicer0   c                    s(   t  | d t | d  | d fS )Nr   r   r   )r#   Zcdiv)ZMETAr'   r(   r	   r
   r      r   z_matmul._call.<locals>.<lambda>r    )r.   r/   )rE   Zstride
contiguousshapetorchemptyr0   float16Zbfloat16float32r3   int32rC   )rA   rB   rE   r)   _cr/   Zgridr	   rF   r
   _callx   s$    

"z_matmul._callc                 C   s   t ||S r   )rD   rP   )ctxrA   rB   r	   r	   r
   forward   s    z_matmul.forwardN)	__name__
__module____qualname__rC   kernelZ_locksstaticmethodrP   rR   r	   r	   r	   r
   rD   s   s   
rD   )rI   r#   Ztriton.languagelanguager3   Zmatmul_perf_modelr   r   r   r&   Zautotuner$   Z
heuristicsZjitZ	constexprrC   ZautogradZFunctionrD   applymatmulr	   r	   r	   r
   <module>   sV   4#