a
    kh                     @   sp  d Z ddlmZ ddlmZ ddlmZmZmZmZm	Z	m
Z
mZmZmZmZmZmZmZmZmZmZmZmZmZmZmZ ddlmZmZmZmZmZm Z m!Z!m"Z"m#Z#m$Z$m%Z%m&Z&m'Z'm(Z(m)Z)m*Z*m+Z+m,Z,m-Z-m.Z.m/Z/m0Z0m1Z1m2Z2m3Z3m4Z4m5Z5m6Z6m7Z7m8Z8m9Z9m:Z:m;Z;m<Z<m=Z=m>Z>m?Z?m@Z@mAZAmBZBmCZCmDZDmEZEmFZFmGZGmHZHmIZImJZJmKZKmLZLmMZMmNZNmOZOmPZPmQZQmRZRmSZSmTZTmUZUmVZVmWZWmXZXmYZYmZZZm[Z[m\Z\m]Z]m^Z^m_Z_m`Z`maZambZbmcZcmdZdmeZemfZfmgZgmhZhmiZimjZjmkZkmlZlmmZmmnZnmoZompZpmqZqmrZr ddlmsZsmtZtmuZumvZvmwZwmxZxmyZymzZzm{Z{m|Z|m}Z}m~Z~mZmZmZmZmZ ddlmZmZmZmZmZmZmZmZmZmZ g dZd	d
 ZdS )zisort:skip_file   )math)extra)argmaxargminbitonic_mergecdivcumprodcumsumflip
interleavemaxminravel	reduce_orsigmoidsoftmaxsortsum	swizzle2dtopkxor_sumzeros
zeros_like)XPropagateNanTRITON_MAX_TENSOR_NUMELload_tensor_descriptorstore_tensor_descriptormake_tensor_descriptortensor_descriptortensor_descriptor_typeaddadvancearangeassociative_scanassume
async_task
atomic_add
atomic_and
atomic_cas
atomic_max
atomic_min	atomic_oratomic_xchg
atomic_xorbfloat16
block_type	broadcastbroadcast_tocatcastclampconst	constexprconstexpr_functiondebug_barrierdevice_assertdevice_printdot
dot_scaleddtypeexpand_dimsfloat16float32float64float8e4b15
float8e4nv
float8e4b8float8e5float8e5b16fullgather	histograminline_asm_elementwiseint1int16int32int64int8joinloadmake_block_ptrmax_constancymax_contiguousmaximumminimummultiple_ofnum_programspermutepi32_tpointer_type
program_idrangereducereshapeslicesplitstatic_assertstatic_printstatic_rangestoretensortranstuple
tuple_typeuint16uint32uint64uint8viewvoidwhere)umulhiexpexp2fmaloglog2cosrsqrtsinsqrtsqrt_rnabsfdivdiv_rnerffloorceil)
pair_uniform_to_normalphiloxphilox_implrandrand4xrandint	randint4xrandnrandn4xuint_to_uniform_float)r   r   r   r   r   r   r|   r    r!   r"   r   r   r#   r$   r%   r&   r'   r(   r)   r*   r+   r,   r-   r.   r   r/   r0   r1   r2   r3   r   r   r4   r5   r6   r7   rw   r   r	   r8   r9   r:   r~   r;   r<   r=   r   rr   rs   r>   r   r}   r
   r?   r@   rA   rB   rC   rD   rE   rF   r   rt   rG   rH   rI   rJ   r   rK   rL   rM   rN   rO   rP   rQ   ru   rv   rR   r   r   rS   rT   rU   r   rV   rW   rX   r   rY   r   r   rZ   r[   r\   r   r   r   r   r   r   r]   r   r^   r   r_   rx   r`   r   ry   r   r   ra   rz   r{   rb   rc   rd   re   r   r   rf   r   rg   rh   rj   rk   rl   rm   r   rq   rn   ro   rp   r   r   r   c                 C   s  ddl m} t| |r<t| jdd }tdd | D |S | d dkr| dd  } d}| d d	krt| dd  } d
}t| }t||dS | 	dr| 
dd d}|j
ddd\}}|j
ddd\}}dd |d
dD }|d}	t|	}
t|}t|}ttg| }ttg| }t||}|
rzddlm} ddlm} t|	t|d}	t|	|slJ |||||	S t|||S | dkrtS ttttttttt t!t"t#ttt!t$t%t&t't!d}||  S )N    )rh   _fieldsc                 S   s   g | ]}t |qS  )	str_to_ty).0xr   r   F/var/www/auris/lib/python3.9/site-packages/triton/language/__init__.py
<listcomp>      zstr_to_ty.<locals>.<listcomp>*r   FkT)Z
element_tyr5   Z
tensordesc<>[   )maxsplit]c                 S   s   g | ]}t | qS r   )intstrip)r   sr   r   r   r   &  r   ,)NVMMASharedLayout)r   r6   )Zfp8e4nvZfp8e4b8Zfp8e5Zfp8e5b16Zfp8e4b15Zfp16Zbf16Zfp32Zfp64i1i8Zi16Zi32Zi64u1u8u16u32Zu64B)(builtinsrh   
isinstancetype__dict__getri   r   r[   
startswithra   rstriplstriplenrM   rN   r/   Z+triton.experimental.gluon.language._layoutsr   Z4triton.experimental.gluon.language.nvidia.hopper.tmar   evaldictr6   rC   rD   rE   rF   rB   r?   r.   r@   rA   rK   rO   rL   rm   rj   rk   rl   )namerh   fieldsr5   tyinnerr=   restZblock_shapeZlayoutZis_gluonndimZ
shape_typeZstride_typeblockr   Zgluon_tensor_descriptor_typeZtysr   r   r   r     sn    



r   N)__doc__ r   r   standardr   r   r   r   r   r	   r
   r   r   r   r   r   r   r   r   r   r   r   r   r   r   corer   r   r   r   r   r   r   r    r!   r"   r#   r$   r%   r&   r'   r(   r)   r*   r+   r,   r-   r.   r/   r0   r1   r2   r3   r4   r5   r6   r7   r8   r9   r:   r;   r<   r=   r>   r?   r@   rA   rB   rC   rD   rE   rF   rG   rH   rI   rJ   rK   rL   rM   rN   rO   rP   rQ   rR   rS   rT   rU   rV   rW   rX   rY   rZ   r[   r\   r]   r^   r_   r`   ra   rb   rc   rd   re   rf   rg   rh   ri   rj   rk   rl   rm   rn   ro   rp   rq   rr   rs   rt   ru   rv   rw   rx   ry   rz   r{   r|   r}   r~   r   r   r   randomr   r   r   r   r   r   r   r   r   r   __all__r   r   r   r   r   <module>   s   \ iZL0 