
    Sܶi>                         d dl mZ d dlZd dlZ	 d dlZd dlmZ n# e	$ r  e
d          w xY wej        dej        fd            Z ed          defd            Zd	ej        defd
ZdS )    )	lru_cacheNz4triton import failed; try `pip install --pre triton`
BLOCK_SIZEc	                 b   t          j        d|          }	|	|k     }
t          d||z   dz             D ]{}t          j                     | |dz
  |z  z   }| ||z  z   }| ||z  z   dz   }t          j        ||	z   |
          }t          j        ||	z   |
          }t          j        ||	z   |
          }t          j        ||dz
  |z  z   |	z   |
d          }|t          j        t          j        ||          |          z   }| |dz   |z  z   dz   }t          j        ||	z   ||
           ||dz   |z  z   dz   }t          j        ||	z   d|
||k    z  ||k    z             t          j        ||	z   d|
||k    z  ||k    z             t          j        ||	z   d|
||k    z  ||k    z             }d S )Nr      mask)r   other   )tlarangerangedebug_barrierloadminimumstore)costtracexx_stridecost_stridetrace_strideNMr   offsetsr   kp0p1p2c0c1c2x_rowcost_rowcost_ptr	trace_ptrs                         U/root/projects/openclaw-proxy/venv/lib/python3.11/site-packages/whisper/triton_ops.py
dtw_kernelr'      s    i:&&GQ;D1a!eai   N N
QUk))AO#AO#a'WR'\---WR'\---WR'\---QUh..8t1MMM2:bjR&8&8"===1q5K//!3
G#XD9999QUl22Q6	
W$adbBh.?28.LMMMM
W$adbBh.?28.LMMMM
W$adbBh.?28.LMMMMM)N N    )maxsizefilter_widthc           
      <    t           j        dt          j        fd            }t          j        |j                  }|j                            dd                    d t                     D                                 }|                    dd                     fdt           d	z  d
z             D                                 }|                    dd d	z             }t          |d          du r|                    |           d |_        n||_        |S )Nr   c                     t          j        d          }t          j        d|          }||k     }|||z  z   }| ||z  z   }	t           t           t          j        |	|z   t          |           d S )Nr   r   )r   
program_idr   LOAD_ALL_ROWS_HEREBUBBLESORT_HEREr   MIDDLE_ROW_HERE)
yr   r   y_strider   row_idxr   r   x_ptry_ptrs
             r&   kernelzmedian_kernel.<locals>.kernel-   sx     -"")Az**!Gh&&Gh&&
/======r(   z    LOAD_ALL_ROWS_HERE
c                      g | ]}d | d| dS )    rowz = tl.load(x_ptr + offsets + z, mask=mask) ).0is     r&   
<listcomp>z!median_kernel.<locals>.<listcomp>B   s<        J!II!III  r(   z    BUBBLESORT_HERE

c           	      t    g | ]4}d                      d t          |z
  dz
            D                       5S )r>   c                     g | ]L}d                      d| d|dz    d| d|dz    d	d| d|dz    d| d|dz    d	d| d	d|dz    d
g          MS )r7   z    smaller = tl.where(rowz < rowr   z, row)z    larger = tl.where(rowz > rowr9   z
 = smallerz	 = larger)join)r;   js     r&   r=   z,median_kernel.<locals>.<listcomp>.<listcomp>N   s     
 
 
  		 bQ b ba!e b bRS b bZ[^_Z_ b b b aA a aQU a aQR a aYZ]^Y^ a a a 7! 7 7 7 :!a% : : :	 
 
 
r(   r   )rB   r   )r;   r<   r*   s     r&   r=   z!median_kernel.<locals>.<listcomp>L   si         
 
 "'|a'7!';!<!<
 
 
   r(   r
   r   r0   row_unsafe_update_srcT)tritonjitr   	constexprJITFunctionfnsrcreplacerB   r   hasattrrE   hash)r*   r6   
new_kernels   `  r&   median_kernelrP   +   sZ   Z>.0l> > > Z>  	**F## 		 |,,  	
 	
 J ##    |q01455  	
 	
 J, ##$57P\Q=N7P7PQQJv+,,44!!*---
Mr(   r   c                    |                                                      d|d          }t          j        |j        dd                   }t          |          }t          j        |d                   }d|                    d          dz
  	                                z  } ||f         || |                     d          |                    d          |           |S )zBApply a median filter of given width along the last dimension of xr   N).r   )r   )

contiguousunfoldnpprodshaperP   torch
empty_likestride
bit_length)r   r*   slicesgridr6   r1   r   s          r&   median_filter_cudar_   j   s    \\^^""2|Q77F76<$%%D<((F((Aqxx||a'33555JFD7OAq!((2,,LLLLHr(   )	functoolsr   numpyrV   rY   rF   triton.languagelanguager   ImportErrorRuntimeErrorrG   rH   r'   intrP   Tensorr_   r:   r(   r&   <module>rh      s             OMMM        O O O
,M
N
NNO NKM<N N N N: 4; ; ; ; ;|%, c      s   
 ,